diff --git a/BUILD.gn b/BUILD.gn index a37e2a0629..f54e12a083 100644 --- a/BUILD.gn +++ b/BUILD.gn @@ -900,6 +900,8 @@ source_set("tint_unittests_wgsl_reader_src") { "src/reader/wgsl/parser_impl_exclusive_or_expression_test.cc", "src/reader/wgsl/parser_impl_for_stmt_test.cc", "src/reader/wgsl/parser_impl_function_decl_test.cc", + "src/reader/wgsl/parser_impl_function_decoration_list_test.cc", + "src/reader/wgsl/parser_impl_function_decoration_test.cc", "src/reader/wgsl/parser_impl_function_header_test.cc", "src/reader/wgsl/parser_impl_function_type_decl_test.cc", "src/reader/wgsl/parser_impl_global_constant_decl_test.cc", diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 04fbea405d..c87bd64ea6 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -431,6 +431,8 @@ if(${TINT_BUILD_WGSL_READER}) reader/wgsl/parser_impl_exclusive_or_expression_test.cc reader/wgsl/parser_impl_for_stmt_test.cc reader/wgsl/parser_impl_function_decl_test.cc + reader/wgsl/parser_impl_function_decoration_list_test.cc + reader/wgsl/parser_impl_function_decoration_test.cc reader/wgsl/parser_impl_function_header_test.cc reader/wgsl/parser_impl_function_type_decl_test.cc reader/wgsl/parser_impl_global_constant_decl_test.cc diff --git a/src/ast/builtin.cc b/src/ast/builtin.cc index b0844cf771..d5818a14e8 100644 --- a/src/ast/builtin.cc +++ b/src/ast/builtin.cc @@ -47,10 +47,6 @@ std::ostream& operator<<(std::ostream& out, Builtin builtin) { out << "frag_depth"; break; } - case Builtin::kWorkgroupSize: { - out << "workgroup_size"; - break; - } case Builtin::kLocalInvocationId: { out << "local_invocation_id"; break; diff --git a/src/ast/builtin.h b/src/ast/builtin.h index 7fb3fe5af1..5d0d763e66 100644 --- a/src/ast/builtin.h +++ b/src/ast/builtin.h @@ -29,7 +29,6 @@ enum class Builtin { kFrontFacing, kFragCoord, kFragDepth, - kWorkgroupSize, kLocalInvocationId, kLocalInvocationIdx, kGlobalInvocationId diff --git a/src/ast/function.h b/src/ast/function.h index d64c0236b8..078f69293b 100644 --- a/src/ast/function.h +++ b/src/ast/function.h @@ -82,6 +82,12 @@ class Function : public Node { /// @returns the function params const VariableList& params() const { return params_; } + /// Sets the function decorations + /// @param decos the decorations to set. This will overwrite any existing + /// decorations + void set_decorations(ast::FunctionDecorationList decos) { + decorations_ = std::move(decos); + } /// Adds a decoration to the function /// @param deco the decoration to set void add_decoration(std::unique_ptr deco) { diff --git a/src/reader/spirv/enum_converter.cc b/src/reader/spirv/enum_converter.cc index 2b54362305..1efc52354f 100644 --- a/src/reader/spirv/enum_converter.cc +++ b/src/reader/spirv/enum_converter.cc @@ -80,8 +80,6 @@ ast::Builtin EnumConverter::ToBuiltin(SpvBuiltIn b) { return ast::Builtin::kFragCoord; case SpvBuiltInFragDepth: return ast::Builtin::kFragDepth; - case SpvBuiltInWorkgroupSize: - return ast::Builtin::kWorkgroupSize; case SpvBuiltInLocalInvocationId: return ast::Builtin::kLocalInvocationId; case SpvBuiltInLocalInvocationIndex: diff --git a/src/reader/spirv/enum_converter_test.cc b/src/reader/spirv/enum_converter_test.cc index fb9c8d0f04..6d542255cd 100644 --- a/src/reader/spirv/enum_converter_test.cc +++ b/src/reader/spirv/enum_converter_test.cc @@ -215,8 +215,6 @@ INSTANTIATE_TEST_SUITE_P( BuiltinCase{SpvBuiltInFrontFacing, true, ast::Builtin::kFrontFacing}, BuiltinCase{SpvBuiltInFragCoord, true, ast::Builtin::kFragCoord}, BuiltinCase{SpvBuiltInFragDepth, true, ast::Builtin::kFragDepth}, - BuiltinCase{SpvBuiltInWorkgroupSize, true, - ast::Builtin::kWorkgroupSize}, BuiltinCase{SpvBuiltInLocalInvocationId, true, ast::Builtin::kLocalInvocationId}, BuiltinCase{SpvBuiltInLocalInvocationIndex, true, diff --git a/src/reader/wgsl/lexer.cc b/src/reader/wgsl/lexer.cc index 37bfeb5238..4d2a1849fa 100644 --- a/src/reader/wgsl/lexer.cc +++ b/src/reader/wgsl/lexer.cc @@ -729,7 +729,8 @@ Token Lexer::check_keyword(const Source& source, const std::string& str) { return {Token::Type::kVoid, source, "void"}; if (str == "workgroup") return {Token::Type::kWorkgroup, source, "workgroup"}; - + if (str == "workgroup_size") + return {Token::Type::kWorkgroupSize, source, "workgroup_size"}; return {}; } diff --git a/src/reader/wgsl/lexer_test.cc b/src/reader/wgsl/lexer_test.cc index d21f3b66e3..b722246ff7 100644 --- a/src/reader/wgsl/lexer_test.cc +++ b/src/reader/wgsl/lexer_test.cc @@ -542,7 +542,8 @@ INSTANTIATE_TEST_SUITE_P( TokenData{"vec4", Token::Type::kVec4}, TokenData{"vertex", Token::Type::kVertex}, TokenData{"void", Token::Type::kVoid}, - TokenData{"workgroup", Token::Type::kWorkgroup})); + TokenData{"workgroup", Token::Type::kWorkgroup}, + TokenData{"workgroup_size", Token::Type::kWorkgroupSize})); using KeywordTest_Reserved = testing::TestWithParam; TEST_P(KeywordTest_Reserved, Parses) { diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc index ae83f80dd3..c3f79f3dcf 100644 --- a/src/reader/wgsl/parser_impl.cc +++ b/src/reader/wgsl/parser_impl.cc @@ -63,6 +63,7 @@ #include "src/ast/unary_op.h" #include "src/ast/unary_op_expression.h" #include "src/ast/variable_decl_statement.h" +#include "src/ast/workgroup_decoration.h" #include "src/reader/wgsl/lexer.h" #include "src/type_manager.h" @@ -95,9 +96,6 @@ ast::Builtin ident_to_builtin(const std::string& str) { if (str == "frag_depth") { return ast::Builtin::kFragDepth; } - if (str == "workgroup_size") { - return ast::Builtin::kWorkgroupSize; - } if (str == "local_invocation_id") { return ast::Builtin::kLocalInvocationId; } @@ -110,6 +108,14 @@ ast::Builtin ident_to_builtin(const std::string& str) { return ast::Builtin::kNone; } +bool IsVariableDecoration(Token t) { + return t.IsLocation() || t.IsBuiltin() || t.IsBinding() || t.IsSet(); +} + +bool IsFunctionDecoration(Token t) { + return t.IsWorkgroupSize(); +} + } // namespace ParserImpl::ParserImpl(Context* ctx, const std::string& input) @@ -444,18 +450,25 @@ ast::VariableDecorationList ParserImpl::variable_decoration_list() { if (!t.IsAttrLeft()) return decos; + // Check the empty list before verifying the contents + t = peek(1); + if (t.IsAttrRight()) { + set_error(t, "empty variable decoration list"); + return {}; + } + + // Make sure we're looking at variable decorations not some other kind + if (!IsVariableDecoration(peek(1))) { + return decos; + } + next(); // consume the peek auto deco = variable_decoration(); if (has_error()) return {}; if (deco == nullptr) { - t = peek(); - if (t.IsAttrRight()) { - set_error(t, "empty variable decoration list"); - return {}; - } - set_error(t, "missing variable decoration for decoration list"); + set_error(peek(), "missing variable decoration for decoration list"); return {}; } for (;;) { @@ -1738,13 +1751,29 @@ ParserImpl::struct_member_decoration() { } // function_decl -// : function_header body_stmt +// : function_decoration_decl* function_header body_stmt std::unique_ptr ParserImpl::function_decl() { + ast::FunctionDecorationList decos; + for (;;) { + size_t s = decos.size(); + if (!function_decoration_decl(decos)) { + return nullptr; + } + if (decos.size() == s) { + break; + } + } + auto f = function_header(); if (has_error()) return nullptr; - if (f == nullptr) + if (f == nullptr) { + if (decos.size() > 0) { + set_error(peek(), "error parsing function declaration"); + } return nullptr; + } + f->set_decorations(std::move(decos)); auto body = body_stmt(); if (has_error()) @@ -1754,6 +1783,131 @@ std::unique_ptr ParserImpl::function_decl() { return f; } +// function_decoration_decl +// : ATTR_LEFT (function_decoration COMMA)* function_decoration ATTR_RIGHT +bool ParserImpl::function_decoration_decl(ast::FunctionDecorationList& decos) { + auto t = peek(); + if (!t.IsAttrLeft()) { + return true; + } + // Handle error on empty attributes before the type check + t = peek(1); + if (t.IsAttrRight()) { + set_error(t, "missing decorations for function decoration block"); + return false; + } + + // Make sure we're looking at function decorations and not some other kind + if (!IsFunctionDecoration(peek(1))) { + return true; + } + + next(); // Consume the peek + + size_t count = 0; + for (;;) { + auto deco = function_decoration(); + if (has_error()) { + return false; + } + if (deco == nullptr) { + set_error(peek(), "expected decoration but none found"); + return false; + } + decos.push_back(std::move(deco)); + count++; + + t = peek(); + if (!t.IsComma()) { + break; + } + next(); // Consume the peek + } + if (count == 0) { + set_error(peek(), "missing decorations for function decoration block"); + return false; + } + + t = next(); + if (!t.IsAttrRight()) { + set_error(t, "missing ]] for function decorations"); + return false; + } + return true; +} + +// function_decoration +// : TODO(dsinclair) STAGE PAREN_LEFT pipeline_stage PAREN_RIGHT +// | WORKGROUP_SIZE PAREN_LEFT INT_LITERAL +// (COMMA INT_LITERAL (COMMA INT_LITERAL)?)? PAREN_RIGHT +std::unique_ptr ParserImpl::function_decoration() { + auto t = peek(); + if (t.IsWorkgroupSize()) { + next(); // Consume the peek + + t = next(); + if (!t.IsParenLeft()) { + set_error(t, "missing ( for workgroup_size"); + return nullptr; + } + + t = next(); + if (!t.IsSintLiteral()) { + set_error(t, "missing x value for workgroup_size"); + return nullptr; + } + if (t.to_i32() <= 0) { + set_error(t, "invalid value for workgroup_size x parameter"); + return nullptr; + } + int32_t x = t.to_i32(); + int32_t y = 1; + int32_t z = 1; + + t = peek(); + if (t.IsComma()) { + next(); // Consume the peek + + t = next(); + if (!t.IsSintLiteral()) { + set_error(t, "missing y value for workgroup_size"); + return nullptr; + } + if (t.to_i32() <= 0) { + set_error(t, "invalid value for workgroup_size y parameter"); + return nullptr; + } + y = t.to_i32(); + + t = peek(); + if (t.IsComma()) { + next(); // Consume the peek + + t = next(); + if (!t.IsSintLiteral()) { + set_error(t, "missing z value for workgroup_size"); + return nullptr; + } + if (t.to_i32() <= 0) { + set_error(t, "invalid value for workgroup_size z parameter"); + return nullptr; + } + z = t.to_i32(); + } + } + + t = next(); + if (!t.IsParenRight()) { + set_error(t, "missing ) for workgroup_size"); + return nullptr; + } + + return std::make_unique(uint32_t(x), uint32_t(y), + uint32_t(z)); + } + return nullptr; +} + // function_type_decl // : type_decl // | VOID diff --git a/src/reader/wgsl/parser_impl.h b/src/reader/wgsl/parser_impl.h index 6da84df4a9..578b4fdc7e 100644 --- a/src/reader/wgsl/parser_impl.h +++ b/src/reader/wgsl/parser_impl.h @@ -182,6 +182,13 @@ class ParserImpl { /// Parses a `function_decl` grammar element /// @returns the parsed function, nullptr otherwise std::unique_ptr function_decl(); + /// Parses a `function_decoration_decl` grammar element + /// @param decos list to store the parsed decorations + /// @returns true on successful parse; false otherwise + bool function_decoration_decl(ast::FunctionDecorationList& decos); + /// Parses a `function_decoration` grammar element + /// @returns the parsed decoration, nullptr otherwise + std::unique_ptr function_decoration(); /// Parses a `texture_sampler_types` grammar element /// @returns the parsed Type or nullptr if none matched. ast::type::Type* texture_sampler_types(); diff --git a/src/reader/wgsl/parser_impl_function_decl_test.cc b/src/reader/wgsl/parser_impl_function_decl_test.cc index fe8e3c7d7e..3ae5905074 100644 --- a/src/reader/wgsl/parser_impl_function_decl_test.cc +++ b/src/reader/wgsl/parser_impl_function_decl_test.cc @@ -15,6 +15,7 @@ #include "gtest/gtest.h" #include "src/ast/function.h" #include "src/ast/type/type.h" +#include "src/ast/workgroup_decoration.h" #include "src/reader/wgsl/parser_impl.h" #include "src/reader/wgsl/parser_impl_test_helper.h" @@ -45,6 +46,113 @@ TEST_F(ParserImplTest, FunctionDecl) { EXPECT_TRUE(body->get(0)->IsReturn()); } +TEST_F(ParserImplTest, FunctionDecl_DecorationList) { + auto* p = parser("[[workgroup_size(2, 3, 4)]] fn main() -> void { return; }"); + auto f = p->function_decl(); + ASSERT_FALSE(p->has_error()) << p->error(); + ASSERT_NE(f, nullptr); + + EXPECT_EQ(f->name(), "main"); + ASSERT_NE(f->return_type(), nullptr); + EXPECT_TRUE(f->return_type()->IsVoid()); + ASSERT_EQ(f->params().size(), 0u); + ASSERT_NE(f->return_type(), nullptr); + EXPECT_TRUE(f->return_type()->IsVoid()); + + auto& decos = f->decorations(); + ASSERT_EQ(decos.size(), 1u); + ASSERT_TRUE(decos[0]->IsWorkgroup()); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + std::tie(x, y, z) = decos[0]->AsWorkgroup()->values(); + EXPECT_EQ(x, 2u); + EXPECT_EQ(y, 3u); + EXPECT_EQ(z, 4u); + + auto* body = f->body(); + ASSERT_EQ(body->size(), 1u); + EXPECT_TRUE(body->get(0)->IsReturn()); +} + +TEST_F(ParserImplTest, FunctionDecl_DecorationList_MultipleEntries) { + auto* p = parser(R"( +[[workgroup_size(2, 3, 4), workgroup_size(5, 6, 7)]] +fn main() -> void { return; })"); + auto f = p->function_decl(); + ASSERT_FALSE(p->has_error()) << p->error(); + ASSERT_NE(f, nullptr); + + EXPECT_EQ(f->name(), "main"); + ASSERT_NE(f->return_type(), nullptr); + EXPECT_TRUE(f->return_type()->IsVoid()); + ASSERT_EQ(f->params().size(), 0u); + ASSERT_NE(f->return_type(), nullptr); + EXPECT_TRUE(f->return_type()->IsVoid()); + + auto& decos = f->decorations(); + ASSERT_EQ(decos.size(), 2u); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + ASSERT_TRUE(decos[0]->IsWorkgroup()); + std::tie(x, y, z) = decos[0]->AsWorkgroup()->values(); + EXPECT_EQ(x, 2u); + EXPECT_EQ(y, 3u); + EXPECT_EQ(z, 4u); + + ASSERT_TRUE(decos[1]->IsWorkgroup()); + std::tie(x, y, z) = decos[1]->AsWorkgroup()->values(); + EXPECT_EQ(x, 5u); + EXPECT_EQ(y, 6u); + EXPECT_EQ(z, 7u); + + auto* body = f->body(); + ASSERT_EQ(body->size(), 1u); + EXPECT_TRUE(body->get(0)->IsReturn()); +} + +TEST_F(ParserImplTest, FunctionDecl_DecorationList_MultipleLists) { + auto* p = parser(R"( +[[workgroup_size(2, 3, 4)]] +[[workgroup_size(5, 6, 7)]] +fn main() -> void { return; })"); + auto f = p->function_decl(); + ASSERT_FALSE(p->has_error()) << p->error(); + ASSERT_NE(f, nullptr); + + EXPECT_EQ(f->name(), "main"); + ASSERT_NE(f->return_type(), nullptr); + EXPECT_TRUE(f->return_type()->IsVoid()); + ASSERT_EQ(f->params().size(), 0u); + ASSERT_NE(f->return_type(), nullptr); + EXPECT_TRUE(f->return_type()->IsVoid()); + + auto& decos = f->decorations(); + ASSERT_EQ(decos.size(), 2u); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + ASSERT_TRUE(decos[0]->IsWorkgroup()); + std::tie(x, y, z) = decos[0]->AsWorkgroup()->values(); + EXPECT_EQ(x, 2u); + EXPECT_EQ(y, 3u); + EXPECT_EQ(z, 4u); + + ASSERT_TRUE(decos[1]->IsWorkgroup()); + std::tie(x, y, z) = decos[1]->AsWorkgroup()->values(); + EXPECT_EQ(x, 5u); + EXPECT_EQ(y, 6u); + EXPECT_EQ(z, 7u); + + auto* body = f->body(); + ASSERT_EQ(body->size(), 1u); + EXPECT_TRUE(body->get(0)->IsReturn()); +} + TEST_F(ParserImplTest, FunctionDecl_InvalidHeader) { auto* p = parser("fn main() -> { }"); auto f = p->function_decl(); diff --git a/src/reader/wgsl/parser_impl_function_decoration_list_test.cc b/src/reader/wgsl/parser_impl_function_decoration_list_test.cc new file mode 100644 index 0000000000..9717eed5c0 --- /dev/null +++ b/src/reader/wgsl/parser_impl_function_decoration_list_test.cc @@ -0,0 +1,98 @@ +// Copyright 2020 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "gtest/gtest.h" +#include "src/ast/workgroup_decoration.h" +#include "src/reader/wgsl/parser_impl.h" +#include "src/reader/wgsl/parser_impl_test_helper.h" + +namespace tint { +namespace reader { +namespace wgsl { +namespace { + +TEST_F(ParserImplTest, FunctionDecorationList_Parses) { + auto* p = parser("[[workgroup_size(2), workgroup_size(3, 4, 5)]]"); + ast::FunctionDecorationList decos; + ASSERT_TRUE(p->function_decoration_decl(decos)); + ASSERT_FALSE(p->has_error()) << p->error(); + ASSERT_EQ(decos.size(), 2u); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + ASSERT_TRUE(decos[0]->IsWorkgroup()); + std::tie(x, y, z) = decos[0]->AsWorkgroup()->values(); + EXPECT_EQ(x, 2u); + + ASSERT_TRUE(decos[1]->IsWorkgroup()); + std::tie(x, y, z) = decos[1]->AsWorkgroup()->values(); + EXPECT_EQ(x, 3u); + EXPECT_EQ(y, 4u); + EXPECT_EQ(z, 5u); +} + +TEST_F(ParserImplTest, FunctionDecorationList_Empty) { + auto* p = parser("[[]]"); + ast::FunctionDecorationList decos; + ASSERT_FALSE(p->function_decoration_decl(decos)); + ASSERT_TRUE(p->has_error()); + ASSERT_EQ(p->error(), + "1:3: missing decorations for function decoration block"); +} + +TEST_F(ParserImplTest, FunctionDecorationList_Invalid) { + auto* p = parser("[[invalid]]"); + ast::FunctionDecorationList decos; + ASSERT_TRUE(p->function_decoration_decl(decos)); + ASSERT_FALSE(p->has_error()); + ASSERT_TRUE(decos.empty()); +} + +TEST_F(ParserImplTest, FunctionDecorationList_ExtraComma) { + auto* p = parser("[[workgroup_size(2), ]]"); + ast::FunctionDecorationList decos; + ASSERT_FALSE(p->function_decoration_decl(decos)); + ASSERT_TRUE(p->has_error()); + ASSERT_EQ(p->error(), "1:22: expected decoration but none found"); +} + +TEST_F(ParserImplTest, FunctionDecorationList_MissingComma) { + auto* p = parser("[[workgroup_size(2) workgroup_size(2)]]"); + ast::FunctionDecorationList decos; + ASSERT_FALSE(p->function_decoration_decl(decos)); + ASSERT_TRUE(p->has_error()); + ASSERT_EQ(p->error(), "1:21: missing ]] for function decorations"); +} + +TEST_F(ParserImplTest, FunctionDecorationList_BadDecoration) { + auto* p = parser("[[workgroup_size()]]"); + ast::FunctionDecorationList decos; + ASSERT_FALSE(p->function_decoration_decl(decos)); + ASSERT_TRUE(p->has_error()); + ASSERT_EQ(p->error(), "1:18: missing x value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecorationList_MissingRightAttr) { + auto* p = parser("[[workgroup_size(2), workgroup_size(3, 4, 5)"); + ast::FunctionDecorationList decos; + ASSERT_FALSE(p->function_decoration_decl(decos)); + ASSERT_TRUE(p->has_error()); + ASSERT_EQ(p->error(), "1:45: missing ]] for function decorations"); +} + +} // namespace +} // namespace wgsl +} // namespace reader +} // namespace tint diff --git a/src/reader/wgsl/parser_impl_function_decoration_test.cc b/src/reader/wgsl/parser_impl_function_decoration_test.cc new file mode 100644 index 0000000000..c7acb60ded --- /dev/null +++ b/src/reader/wgsl/parser_impl_function_decoration_test.cc @@ -0,0 +1,196 @@ +// Copyright 2020 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "gtest/gtest.h" +#include "src/ast/workgroup_decoration.h" +#include "src/reader/wgsl/parser_impl.h" +#include "src/reader/wgsl/parser_impl_test_helper.h" + +namespace tint { +namespace reader { +namespace wgsl { +namespace { + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup) { + auto* p = parser("workgroup_size(4)"); + auto deco = p->function_decoration(); + ASSERT_NE(deco, nullptr); + ASSERT_FALSE(p->has_error()); + ASSERT_TRUE(deco->IsWorkgroup()); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + std::tie(x, y, z) = deco->AsWorkgroup()->values(); + EXPECT_EQ(x, 4u); + EXPECT_EQ(y, 1u); + EXPECT_EQ(z, 1u); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_2Param) { + auto* p = parser("workgroup_size(4, 5)"); + auto deco = p->function_decoration(); + ASSERT_NE(deco, nullptr) << p->error(); + ASSERT_FALSE(p->has_error()); + ASSERT_TRUE(deco->IsWorkgroup()); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + std::tie(x, y, z) = deco->AsWorkgroup()->values(); + EXPECT_EQ(x, 4u); + EXPECT_EQ(y, 5u); + EXPECT_EQ(z, 1u); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_3Param) { + auto* p = parser("workgroup_size(4, 5, 6)"); + auto deco = p->function_decoration(); + ASSERT_NE(deco, nullptr); + ASSERT_FALSE(p->has_error()); + ASSERT_TRUE(deco->IsWorkgroup()); + + uint32_t x = 0; + uint32_t y = 0; + uint32_t z = 0; + std::tie(x, y, z) = deco->AsWorkgroup()->values(); + EXPECT_EQ(x, 4u); + EXPECT_EQ(y, 5u); + EXPECT_EQ(z, 6u); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_TooManyValues) { + auto* p = parser("workgroup_size(1, 2, 3, 4)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:23: missing ) for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Invalid_X_Value) { + auto* p = parser("workgroup_size(-2, 5, 6)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:16: invalid value for workgroup_size x parameter"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Invalid_Y_Value) { + auto* p = parser("workgroup_size(4, 0, 6)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:19: invalid value for workgroup_size y parameter"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Invalid_Z_Value) { + auto* p = parser("workgroup_size(4, 5, -3)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:22: invalid value for workgroup_size z parameter"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_MissingLeftParam) { + auto* p = parser("workgroup_size 4, 5, 6)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:16: missing ( for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_MissingRightParam) { + auto* p = parser("workgroup_size(4, 5, 6"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:23: missing ) for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_MissingValues) { + auto* p = parser("workgroup_size()"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:16: missing x value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_X_Value) { + auto* p = parser("workgroup_size(, 2, 3)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:16: missing x value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_Y_Comma) { + auto* p = parser("workgroup_size(1 2, 3)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:18: missing ) for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_Y_Value) { + auto* p = parser("workgroup_size(1, , 3)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:19: missing y value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_Z_Comma) { + auto* p = parser("workgroup_size(1, 2 3)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:21: missing ) for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_Z_Value) { + auto* p = parser("workgroup_size(1, 2, )"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:22: missing z value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_X_Invalid) { + auto* p = parser("workgroup_size(nan)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:16: missing x value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_Y_Invalid) { + auto* p = parser("workgroup_size(2, nan)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:19: missing y value for workgroup_size"); +} + +TEST_F(ParserImplTest, FunctionDecoration_Workgroup_Missing_Z_Invalid) { + auto* p = parser("workgroup_size(2, 3, nan)"); + auto deco = p->function_decoration(); + ASSERT_EQ(deco, nullptr); + ASSERT_TRUE(p->has_error()); + EXPECT_EQ(p->error(), "1:22: missing z value for workgroup_size"); +} + +} // namespace +} // namespace wgsl +} // namespace reader +} // namespace tint diff --git a/src/reader/wgsl/parser_impl_global_decl_test.cc b/src/reader/wgsl/parser_impl_global_decl_test.cc index 489f07dde9..c04735b6e7 100644 --- a/src/reader/wgsl/parser_impl_global_decl_test.cc +++ b/src/reader/wgsl/parser_impl_global_decl_test.cc @@ -166,6 +166,16 @@ TEST_F(ParserImplTest, GlobalDecl_Function) { EXPECT_EQ(m.functions()[0]->name(), "main"); } +TEST_F(ParserImplTest, GlobalDecl_Function_WithDecoration) { + auto* p = parser("[[workgroup_size(2)]] fn main() -> void { return; }"); + p->global_decl(); + ASSERT_FALSE(p->has_error()) << p->error(); + + auto m = p->module(); + ASSERT_EQ(m.functions().size(), 1u); + EXPECT_EQ(m.functions()[0]->name(), "main"); +} + TEST_F(ParserImplTest, GlobalDecl_Function_Invalid) { auto* p = parser("fn main() -> { return; }"); p->global_decl(); diff --git a/src/reader/wgsl/parser_impl_variable_decoration_list_test.cc b/src/reader/wgsl/parser_impl_variable_decoration_list_test.cc index 33719cab5a..942abe9f3e 100644 --- a/src/reader/wgsl/parser_impl_variable_decoration_list_test.cc +++ b/src/reader/wgsl/parser_impl_variable_decoration_list_test.cc @@ -44,8 +44,8 @@ TEST_F(ParserImplTest, VariableDecorationList_Empty) { TEST_F(ParserImplTest, VariableDecorationList_Invalid) { auto* p = parser(R"([[invalid]])"); auto decos = p->variable_decoration_list(); - ASSERT_TRUE(p->has_error()); - ASSERT_EQ(p->error(), "1:3: missing variable decoration for decoration list"); + ASSERT_FALSE(p->has_error()); + ASSERT_TRUE(decos.empty()); } TEST_F(ParserImplTest, VariableDecorationList_ExtraComma) { diff --git a/src/reader/wgsl/parser_impl_variable_decoration_test.cc b/src/reader/wgsl/parser_impl_variable_decoration_test.cc index 070e8ec1c0..d410b6dbd0 100644 --- a/src/reader/wgsl/parser_impl_variable_decoration_test.cc +++ b/src/reader/wgsl/parser_impl_variable_decoration_test.cc @@ -117,7 +117,6 @@ INSTANTIATE_TEST_SUITE_P( BuiltinData{"front_facing", ast::Builtin::kFrontFacing}, BuiltinData{"frag_coord", ast::Builtin::kFragCoord}, BuiltinData{"frag_depth", ast::Builtin::kFragDepth}, - BuiltinData{"workgroup_size", ast::Builtin::kWorkgroupSize}, BuiltinData{"local_invocation_id", ast::Builtin::kLocalInvocationId}, BuiltinData{"local_invocation_idx", ast::Builtin::kLocalInvocationIdx}, BuiltinData{"global_invocation_id", diff --git a/src/reader/wgsl/token.cc b/src/reader/wgsl/token.cc index dd1ad3ee30..ed90e8a4af 100644 --- a/src/reader/wgsl/token.cc +++ b/src/reader/wgsl/token.cc @@ -349,6 +349,8 @@ std::string Token::TypeToName(Type type) { return "void"; case Token::Type::kWorkgroup: return "workgroup"; + case Token::Type::kWorkgroupSize: + return "workgroup_size"; } return ""; diff --git a/src/reader/wgsl/token.h b/src/reader/wgsl/token.h index da18df348e..0967cd704a 100644 --- a/src/reader/wgsl/token.h +++ b/src/reader/wgsl/token.h @@ -359,7 +359,9 @@ class Token { /// A 'void' kVoid, /// A 'workgroup' - kWorkgroup + kWorkgroup, + /// A 'workgroup_size' + kWorkgroupSize, }; /// Converts a token type to a name @@ -777,6 +779,8 @@ class Token { bool IsVoid() const { return type_ == Type::kVoid; } /// @returns true if token is a 'workgroup' bool IsWorkgroup() const { return type_ == Type::kWorkgroup; } + /// @returns true if token is a 'workgroup_size' + bool IsWorkgroupSize() const { return type_ == Type::kWorkgroupSize; } /// @returns the source line of the token size_t line() const { return source_.line; } diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index d34e572a20..f81e36c6bf 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -1395,11 +1395,6 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const { return "SV_Position"; case ast::Builtin::kFragDepth: return "SV_Depth"; - // TODO(dsinclair): Ignore for now. This has been removed as a builtin - // in the spec. Need to update Tint to match. - // https://github.com/gpuweb/gpuweb/pull/824 - case ast::Builtin::kWorkgroupSize: - return ""; case ast::Builtin::kLocalInvocationId: return "SV_GroupThreadID"; case ast::Builtin::kLocalInvocationIdx: diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc index 6b2324a41f..39b39dca79 100644 --- a/src/writer/hlsl/generator_impl_test.cc +++ b/src/writer/hlsl/generator_impl_test.cc @@ -88,7 +88,6 @@ INSTANTIATE_TEST_SUITE_P( HlslBuiltinData{ast::Builtin::kFrontFacing, "SV_IsFrontFacing"}, HlslBuiltinData{ast::Builtin::kFragCoord, "SV_Position"}, HlslBuiltinData{ast::Builtin::kFragDepth, "SV_Depth"}, - HlslBuiltinData{ast::Builtin::kWorkgroupSize, ""}, HlslBuiltinData{ast::Builtin::kLocalInvocationId, "SV_GroupThreadID"}, HlslBuiltinData{ast::Builtin::kLocalInvocationIdx, "SV_GroupIndex"}, HlslBuiltinData{ast::Builtin::kGlobalInvocationId, diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 6bc21fd079..2251e3b48b 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -1257,11 +1257,6 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const { return "position"; case ast::Builtin::kFragDepth: return "depth(any)"; - // TODO(dsinclair): Ignore for now. This has been removed as a builtin - // in the spec. Need to update Tint to match. - // https://github.com/gpuweb/gpuweb/pull/824 - case ast::Builtin::kWorkgroupSize: - return ""; case ast::Builtin::kLocalInvocationId: return "thread_position_in_threadgroup"; case ast::Builtin::kLocalInvocationIdx: diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc index 922ca9dc52..f57bf52e18 100644 --- a/src/writer/msl/generator_impl_test.cc +++ b/src/writer/msl/generator_impl_test.cc @@ -116,7 +116,6 @@ INSTANTIATE_TEST_SUITE_P( MslBuiltinData{ast::Builtin::kFrontFacing, "front_facing"}, MslBuiltinData{ast::Builtin::kFragCoord, "position"}, MslBuiltinData{ast::Builtin::kFragDepth, "depth(any)"}, - MslBuiltinData{ast::Builtin::kWorkgroupSize, ""}, MslBuiltinData{ast::Builtin::kLocalInvocationId, "thread_position_in_threadgroup"}, MslBuiltinData{ast::Builtin::kLocalInvocationIdx, diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index 879ded399d..050fd1326d 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -2309,8 +2309,6 @@ SpvBuiltIn Builder::ConvertBuiltin(ast::Builtin builtin) const { return SpvBuiltInFragCoord; case ast::Builtin::kFragDepth: return SpvBuiltInFragDepth; - case ast::Builtin::kWorkgroupSize: - return SpvBuiltInWorkgroupSize; case ast::Builtin::kLocalInvocationId: return SpvBuiltInLocalInvocationId; case ast::Builtin::kLocalInvocationIdx: diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc index 5ab3ce2ae3..1e60c0fcfc 100644 --- a/src/writer/spirv/builder_global_variable_test.cc +++ b/src/writer/spirv/builder_global_variable_test.cc @@ -389,7 +389,6 @@ INSTANTIATE_TEST_SUITE_P( BuiltinData{ast::Builtin::kFrontFacing, SpvBuiltInFrontFacing}, BuiltinData{ast::Builtin::kFragCoord, SpvBuiltInFragCoord}, BuiltinData{ast::Builtin::kFragDepth, SpvBuiltInFragDepth}, - BuiltinData{ast::Builtin::kWorkgroupSize, SpvBuiltInWorkgroupSize}, BuiltinData{ast::Builtin::kLocalInvocationId, SpvBuiltInLocalInvocationId}, BuiltinData{ast::Builtin::kLocalInvocationIdx, diff --git a/test/function.wgsl b/test/function.wgsl index e8062f69b5..b7060fb41e 100644 --- a/test/function.wgsl +++ b/test/function.wgsl @@ -15,3 +15,9 @@ fn main() -> f32 { return ((2. * 3.) - 4.) / 5.; } + +[[workgroup_size(2)]] +fn ep() -> void { + return; +} +entry_point compute = ep;