[OpenACC] Implement 'tile' clause parsing

The 'tile' clause takes a 'size-expr-list', where a 'size-expr' is
either an asterisk or an integral constant expression.  This patch
parsess it as an assignment expression, which we'll check for constness
and type in Sema.

Author: erichkeane

clang/include/clang/Basic/OpenACCKinds.h

@@ -241,6 +241,8 @@ enum class OpenACCClauseKind {
   /// 'async' clause, allowed on Compute, Data, 'update', 'wait', and Combined
   /// constructs.
   Async,
+  /// 'tile' clause, allowed on 'loop' and Combined constructs.
+  Tile,
 
   /// Represents an invalid clause, for the purposes of parsing.
   Invalid,
@@ -366,6 +368,9 @@ inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
   case OpenACCClauseKind::Async:
     return Out << "async";
 
+  case OpenACCClauseKind::Tile:
+    return Out << "tile";
+
   case OpenACCClauseKind::Invalid:
     return Out << "<invalid>";
   }

clang/include/clang/Parse/Parser.h

@@ -3592,6 +3592,10 @@ class Parser : public CodeCompletionHandler {
   /// Parses the 'async-argument', which is an integral value with two
   /// 'special' values that are likely negative (but come from Macros).
   ExprResult ParseOpenACCAsyncArgument();
+  /// Parses the 'size-expr', which is an integral value, or an asterisk.
+  bool ParseOpenACCSizeExpr();
+  /// Parses a comma delimited list of 'size-expr's.
+  bool ParseOpenACCSizeExprList();
 
 private:
   //===--------------------------------------------------------------------===//

clang/lib/Parse/ParseOpenACC.cpp

@@ -124,6 +124,7 @@ OpenACCClauseKind getOpenACCClauseKind(Token Tok) {
       .Case("reduction", OpenACCClauseKind::Reduction)
       .Case("self", OpenACCClauseKind::Self)
       .Case("seq", OpenACCClauseKind::Seq)
+      .Case("tile", OpenACCClauseKind::Tile)
       .Case("use_device", OpenACCClauseKind::UseDevice)
       .Case("vector", OpenACCClauseKind::Vector)
       .Case("vector_length", OpenACCClauseKind::VectorLength)
@@ -494,6 +495,7 @@ ClauseParensKind getClauseParensKind(OpenACCDirectiveKind DirKind,
   case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::DeviceType:
   case OpenACCClauseKind::DType:
+  case OpenACCClauseKind::Tile:
     return ClauseParensKind::Required;
 
   case OpenACCClauseKind::Auto:
@@ -618,6 +620,47 @@ bool Parser::ParseOpenACCDeviceTypeList() {
   return false;
 }
 
+/// OpenACC 3.3 Section 2.9:
+/// size-expr is one of:
+//    *
+//    int-expr
+// Note that this is specified under 'gang-arg-list', but also applies to 'tile'
+// via reference.
+bool Parser::ParseOpenACCSizeExpr() {
+  // FIXME: Ensure these are constant expressions.
+
+  // The size-expr ends up being ambiguous when only looking at the current
+  // token, as it could be a deref of a variable/expression.
+  if (getCurToken().is(tok::star) &&
+      NextToken().isOneOf(tok::comma, tok::r_paren)) {
+    ConsumeToken();
+    return false;
+  }
+
+  return getActions()
+      .CorrectDelayedTyposInExpr(ParseAssignmentExpression())
+      .isInvalid();
+}
+
+bool Parser::ParseOpenACCSizeExprList() {
+  if (ParseOpenACCSizeExpr()) {
+    SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
+              Parser::StopBeforeMatch);
+    return false;
+  }
+
+  while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
+    ExpectAndConsume(tok::comma);
+
+    if (ParseOpenACCSizeExpr()) {
+      SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
+                Parser::StopBeforeMatch);
+      return false;
+    }
+  }
+  return false;
+}
+
 // The OpenACC Clause List is a comma or space-delimited list of clauses (see
 // the comment on ParseOpenACCClauseList).  The concept of a 'clause' doesn't
 // really have its owner grammar and each individual one has its own definition.
@@ -757,6 +800,10 @@ bool Parser::ParseOpenACCClauseParams(OpenACCDirectiveKind DirKind,
         return true;
       }
       break;
+    case OpenACCClauseKind::Tile:
+      if (ParseOpenACCSizeExprList())
+        return true;
+      break;
     default:
       llvm_unreachable("Not a required parens type?");
     }

clang/test/ParserOpenACC/parse-clauses.c

@@ -1049,6 +1049,68 @@ void AsyncArgument() {
 #pragma acc parallel async(acc_async_sync)
 }
 
+void Tile() {
+
+  int* Foo;
+  // expected-error@+2{{expected '('}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile
+  for(;;){}
+  // expected-error@+4{{expected expression}}
+  // expected-error@+3{{expected ')'}}
+  // expected-note@+2{{to match this '('}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(
+  for(;;){}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile()
+  for(;;){}
+  // expected-error@+4{{expected expression}}
+  // expected-error@+3{{expected ')'}}
+  // expected-note@+2{{to match this '('}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(,
+  for(;;){}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(,)
+  for(;;){}
+  // expected-error@+2{{use of undeclared identifier 'invalid'}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(returns_int(), *, invalid, *)
+  for(;;){}
+
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(returns_int() *, Foo, *)
+  for(;;){}
+
+  // expected-error@+2{{indirection requires pointer operand ('int' invalid)}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(* returns_int() , *)
+  for(;;){}
+
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(*)
+  for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(*Foo, *Foo)
+  for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(5)
+  for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(*, 5)
+  for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(5, *)
+  for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc loop tile(5, *, 3, *)
+  for(;;){}
+}
+
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc routine worker, vector, seq, nohost
 void bar();