[wgsl-reader] Add workgroup_size parsing

This CL adds parsing of the `workgroup_size` function decoration.

Change-Id: Ia90efc2c014ac0e1614429280cc903d30cf8171d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/28663
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Sarah Mashayekhi <sarahmashay@google.com>

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",

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

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;

src/ast/builtin.h

@@ -29,7 +29,6 @@ enum class Builtin {
   kFrontFacing,
   kFragCoord,
   kFragDepth,
-  kWorkgroupSize,
   kLocalInvocationId,
   kLocalInvocationIdx,
   kGlobalInvocationId

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<FunctionDecoration> deco) {

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:

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,

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 {};
 }
 

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<const char*>;
 TEST_P(KeywordTest_Reserved, Parses) {

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<ast::Function> 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<ast::Function> 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<ast::FunctionDecoration> 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<ast::WorkgroupDecoration>(uint32_t(x), uint32_t(y),
+                                                      uint32_t(z));
+  }
+  return nullptr;
+}
+
 // function_type_decl
 //   : type_decl
 //   | VOID

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<ast::Function> 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<ast::FunctionDecoration> function_decoration();
   /// Parses a `texture_sampler_types` grammar element
   /// @returns the parsed Type or nullptr if none matched.
   ast::type::Type* texture_sampler_types();

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();

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

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

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();

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) {

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",

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 "<unknown>";

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; }

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:

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,

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:

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,

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:

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,

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;