[OpenACC] Implement 'bind' clause parsing.

'bind' takes either a string literal, or an 'identifier' representing
the device-side function that this routine is intended to 'bind' to
(that is, to call).  However, it seems that the intent is to permit the
'identifier' to reference any function, thus we're implementing this as
an ID expression.

Additionally, while working on this I discovered that the 'routine' ID
expression parsing for C++ wouldn't allow non-static member functions to
be referenced since it expected us to call it, this was fixed as a part
of this patch as the 'bind' support needed it too.  A test was added for
routine.

Author: erichkeane

clang/include/clang/Basic/DiagnosticParseKinds.td

@@ -1374,6 +1374,7 @@ def err_acc_expected_reduction_operator
 def err_acc_invalid_reduction_operator
     : Error<"invalid reduction operator,  expected '+', '*', 'max', 'min', "
             "'&', '|', '^', '&&', or '||'">;
+def err_acc_incorrect_bind_arg : Error<"expected identifier or string literal">;
 
 // OpenMP support.
 def warn_pragma_omp_ignored : Warning<

clang/include/clang/Basic/OpenACCKinds.h

@@ -219,6 +219,8 @@ enum class OpenACCClauseKind {
   Reduction,
   /// 'collapse' clause, allowed on 'loop' and Combined constructs.
   Collapse,
+  /// 'bind' clause, allowed on routine constructs.
+  Bind,
 
   /// Represents an invalid clause, for the purposes of parsing.
   Invalid,
@@ -317,6 +319,9 @@ inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
   case OpenACCClauseKind::Collapse:
     return Out << "collapse";
 
+  case OpenACCClauseKind::Bind:
+    return Out << "bind";
+
   case OpenACCClauseKind::Invalid:
     return Out << "<invalid>";
   }

clang/include/clang/Parse/Parser.h

@@ -3581,6 +3581,9 @@ class Parser : public CodeCompletionHandler {
   /// Parses the clause-list for an OpenACC directive.
   void ParseOpenACCClauseList(OpenACCDirectiveKind DirKind);
   bool ParseOpenACCWaitArgument();
+  /// Parses the clause of the 'bind' argument, which can be a string literal or
+  /// an ID expression.
+  ExprResult ParseOpenACCBindClauseArgument();
 
 private:
   //===--------------------------------------------------------------------===//

clang/lib/Parse/ParseOpenACC.cpp

@@ -91,6 +91,7 @@ OpenACCClauseKind getOpenACCClauseKind(Token Tok) {
              Tok.getIdentifierInfo()->getName())
       .Case("attach", OpenACCClauseKind::Attach)
       .Case("auto", OpenACCClauseKind::Auto)
+      .Case("bind", OpenACCClauseKind::Bind)
       .Case("create", OpenACCClauseKind::Create)
       .Case("collapse", OpenACCClauseKind::Collapse)
       .Case("copy", OpenACCClauseKind::Copy)
@@ -467,6 +468,7 @@ ClauseParensKind getClauseParensKind(OpenACCDirectiveKind DirKind,
   case OpenACCClauseKind::Host:
   case OpenACCClauseKind::Reduction:
   case OpenACCClauseKind::Collapse:
+  case OpenACCClauseKind::Bind:
     return ClauseParensKind::Required;
 
   case OpenACCClauseKind::Auto:
@@ -668,6 +670,12 @@ bool Parser::ParseOpenACCClauseParams(OpenACCDirectiveKind DirKind,
         return true;
       break;
     }
+    case OpenACCClauseKind::Bind: {
+      ExprResult BindArg = ParseOpenACCBindClauseArgument();
+      if (BindArg.isInvalid())
+        return true;
+      break;
+    }
     default:
       llvm_unreachable("Not a required parens type?");
     }
@@ -758,7 +766,7 @@ bool Parser::ParseOpenACCWaitArgument() {
 ExprResult Parser::ParseOpenACCIDExpression() {
   ExprResult Res;
   if (getLangOpts().CPlusPlus) {
-    Res = ParseCXXIdExpression(/*isAddressOfOperand=*/false);
+    Res = ParseCXXIdExpression(/*isAddressOfOperand=*/true);
   } else {
     // There isn't anything quite the same as ParseCXXIdExpression for C, so we
     // need to get the identifier, then call into Sema ourselves.
@@ -785,6 +793,25 @@ ExprResult Parser::ParseOpenACCIDExpression() {
   return getActions().CorrectDelayedTyposInExpr(Res);
 }
 
+ExprResult Parser::ParseOpenACCBindClauseArgument() {
+  // OpenACC 3.3 section 2.15:
+  // The bind clause specifies the name to use when calling the procedure on a
+  // device other than the host. If the name is specified as an identifier, it
+  // is called as if that name were specified in the language being compiled. If
+  // the name is specified as a string, the string is used for the procedure
+  // name unmodified.
+  if (getCurToken().is(tok::r_paren)) {
+    Diag(getCurToken(), diag::err_acc_incorrect_bind_arg);
+    return ExprError();
+  }
+
+  if (tok::isStringLiteral(getCurToken().getKind()))
+    return getActions().CorrectDelayedTyposInExpr(ParseStringLiteralExpression(
+        /*AllowUserDefinedLiteral=*/false, /*Unevaluated=*/true));
+
+  return ParseOpenACCIDExpression();
+}
+
 /// OpenACC 3.3, section 1.6:
 /// In this spec, a 'var' (in italics) is one of the following:
 /// - a variable name (a scalar, array, or compisite variable name)

clang/test/ParserOpenACC/parse-clauses.c

@@ -706,3 +706,26 @@ void bar();
 
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc routine(bar) worker, vector, seq, nohost
+
+
+// Bind Clause Parsing.
+
+  // expected-error@+2{{expected '('}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine bind
+void BCP1();
+
+  // expected-error@+2{{expected identifier or string literal}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(BCP1) bind()
+
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine bind("ReductionClauseParsing")
+void BCP2();
+
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(BCP1) bind(BCP2)
+
+  // expected-error@+2{{use of undeclared identifier 'unknown_thing'}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(BCP1) bind(unknown_thing)

clang/test/ParserOpenACC/parse-clauses.cpp

@@ -18,3 +18,56 @@ struct S {
 void use() {
   templ<7, S>();
 }
+
+namespace NS {
+void NSFunc();
+
+class RecordTy { // #RecTy
+  static constexpr bool Value = false; // #VAL
+  void priv_mem_function(); // #PrivMemFun
+  public:
+  static constexpr bool ValuePub = true;
+  void mem_function();
+};
+template<typename T>
+class TemplTy{};
+void function();
+}
+
+
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::NSFunc)
+  // expected-error@+3{{'RecordTy' does not refer to a value}}
+  // expected-note@#RecTy{{declared here}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::RecordTy)
+  // expected-error@+3{{'Value' is a private member of 'NS::RecordTy'}}
+  // expected-note@#VAL{{implicitly declared private here}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::RecordTy::Value)
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::RecordTy::ValuePub)
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::TemplTy<int>)
+  // expected-error@+2{{no member named 'unknown' in namespace 'NS'}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::unknown<int>)
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::function)
+  // expected-error@+3{{'priv_mem_function' is a private member of 'NS::RecordTy'}}
+  // expected-note@#PrivMemFun{{implicitly declared private here}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::RecordTy::priv_mem_function)
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(NS::RecordTy::mem_function)
+
+  // expected-error@+2{{string literal with user-defined suffix cannot be used here}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind("unknown udl"_UDL)
+
+  // expected-warning@+2{{encoding prefix 'u' on an unevaluated string literal has no effect}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(u"16 bits")
+  // expected-warning@+2{{encoding prefix 'U' on an unevaluated string literal has no effect}}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine(use) bind(U"32 bits")

clang/test/ParserOpenACC/parse-constructs.cpp

@@ -5,6 +5,12 @@ namespace NS {
 
   template<typename T>
   void templ(); // expected-note 2{{declared here}}
+
+  class C { // #CDef
+    void private_mem_func(); // #PrivateMemFunc
+    public:
+    void public_mem_func();
+  };
 }
 
 // expected-error@+2{{use of undeclared identifier 'foo'; did you mean 'NS::foo'?}}
@@ -44,3 +50,14 @@ namespace NS {
 // expected-error@+2 {{expected unqualified-id}}
 // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc routine(int)
+
+// expected-error@+3{{'C' does not refer to a value}}
+// expected-note@#CDef{{declared here}}
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine (NS::C)
+// expected-error@+3{{'private_mem_func' is a private member of 'NS::C'}}
+// expected-note@#PrivateMemFunc{{implicitly declared private here}}
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine (NS::C::private_mem_func)
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc routine (NS::C::public_mem_func)