Implement 'bind' restrictions on routine like we previously did for

gang/worker/vector

Discussed separately, this seems like the right way forward.

Author: erichkeane

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -224,29 +224,23 @@ class SemaOpenACCClauseVisitor {
     llvm_unreachable("didn't return from switch above?");
   }
 
-  // Routine has a pretty complicated set of rules for how device_type and the
-  // gang, worker, vector, and seq clauses work.  So diagnose some of it here.
-  bool CheckValidRoutineGangWorkerVectorSeqNewClause(
-      SemaOpenACC::OpenACCParsedClause &Clause) {
+  // Helper for the 'routine' checks during 'new' clause addition. Precondition
+  // is that we already know the new clause is one of the prohbiited ones.
+  template <typename Pred>
+  bool
+  CheckValidRoutineNewClauseHelper(Pred HasPredicate,
+                                   SemaOpenACC::OpenACCParsedClause &Clause) {
     if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Routine)
       return false;
 
-    if (Clause.getClauseKind() != OpenACCClauseKind::Gang &&
-        Clause.getClauseKind() != OpenACCClauseKind::Vector &&
-        Clause.getClauseKind() != OpenACCClauseKind::Worker &&
-        Clause.getClauseKind() != OpenACCClauseKind::Seq)
-      return false;
-    auto ProhibitedPred = llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
-                                        OpenACCVectorClause, OpenACCSeqClause>;
-
     auto *FirstDeviceType =
         llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCDeviceTypeClause>);
 
     if (FirstDeviceType == ExistingClauses.end()) {
       // If there isn't a device type yet, ANY duplicate is wrong.
 
       auto *ExistingProhibitedClause =
-          llvm::find_if(ExistingClauses, ProhibitedPred);
+          llvm::find_if(ExistingClauses, HasPredicate);
 
       if (ExistingProhibitedClause == ExistingClauses.end())
         return false;
@@ -265,8 +259,7 @@ class SemaOpenACCClauseVisitor {
     // between this and the previous 'device_type'.
 
     auto *BeforeDeviceType =
-        std::find_if(ExistingClauses.begin(), FirstDeviceType, ProhibitedPred);
-
+        std::find_if(ExistingClauses.begin(), FirstDeviceType, HasPredicate);
     // If there is one before the device_type (and we know we are after a
     // device_type), than this is ill-formed.
     if (BeforeDeviceType != FirstDeviceType) {
@@ -294,7 +287,7 @@ class SemaOpenACCClauseVisitor {
     auto *LastDeviceType = LastDeviceTypeItr.base() - 1;
 
     auto *ExistingProhibitedSinceLastDevice =
-        std::find_if(LastDeviceType, ExistingClauses.end(), ProhibitedPred);
+        std::find_if(LastDeviceType, ExistingClauses.end(), HasPredicate);
 
     // No prohibited ones since the last device-type.
     if (ExistingProhibitedSinceLastDevice == ExistingClauses.end())
@@ -311,6 +304,35 @@ class SemaOpenACCClauseVisitor {
     return true;
   }
 
+  // Routine has a pretty complicated set of rules for how device_type and the
+  // gang, worker, vector, and seq clauses work.  So diagnose some of it here.
+  bool CheckValidRoutineGangWorkerVectorSeqNewClause(
+      SemaOpenACC::OpenACCParsedClause &Clause) {
+
+    if (Clause.getClauseKind() != OpenACCClauseKind::Gang &&
+        Clause.getClauseKind() != OpenACCClauseKind::Vector &&
+        Clause.getClauseKind() != OpenACCClauseKind::Worker &&
+        Clause.getClauseKind() != OpenACCClauseKind::Seq)
+      return false;
+    auto ProhibitedPred = llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
+                                        OpenACCVectorClause, OpenACCSeqClause>;
+
+    return CheckValidRoutineNewClauseHelper(ProhibitedPred, Clause);
+  }
+
+  // Bind should have similar rules on a routine as gang/worker/vector/seq,
+  // except there is no 'must have 1' rule, so we can get all the checking done
+  // here.
+  bool
+  CheckValidRoutineBindNewClause(SemaOpenACC::OpenACCParsedClause &Clause) {
+
+    if (Clause.getClauseKind() != OpenACCClauseKind::Bind)
+      return false;
+
+    auto HasBindPred = llvm::IsaPred<OpenACCBindClause>;
+    return CheckValidRoutineNewClauseHelper(HasBindPred, Clause);
+  }
+
   // For 'tile' and 'collapse', only allow 1 per 'device_type'.
   template <typename TheClauseTy>
   bool DisallowSinceLastDeviceType(SemaOpenACC::OpenACCParsedClause &Clause) {
@@ -352,7 +374,8 @@ class SemaOpenACCClauseVisitor {
                                          Clause.getClauseKind(),
                                          Clause.getBeginLoc(), ExistingClauses))
       return nullptr;
-    if (CheckValidRoutineGangWorkerVectorSeqNewClause(Clause))
+    if (CheckValidRoutineGangWorkerVectorSeqNewClause(Clause) ||
+        CheckValidRoutineBindNewClause(Clause))
       return nullptr;
 
     switch (Clause.getClauseKind()) {

clang/test/SemaOpenACC/routine-construct-clauses.cpp

@@ -7,11 +7,11 @@ void Func2();
 #pragma acc routine(Func) vector nohost
 #pragma acc routine(Func) nohost seq
 #pragma acc routine(Func) gang
-// expected-error@+2{{OpenACC 'bind' clause cannot appear more than once on a 'routine' directive}}
+// expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}}
 // expected-note@+1{{previous clause is here}}
 #pragma acc routine(Func) gang bind(a) bind(a)
 
-// expected-error@+2{{OpenACC 'bind' clause cannot appear more than once on a 'routine' directive}}
+// expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}}
 // expected-note@+1{{previous clause is here}}
 #pragma acc routine gang bind(a) bind(a)
 void DupeImplName();
@@ -733,3 +733,42 @@ namespace FigureDupesAllowedAroundDeviceType {
   void Func113();
 
 }
+
+namespace BindDupes {
+  // expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq bind(asdf) bind(asdf)
+  void Func1();
+  // expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq bind(asdf) bind(asdf) device_type(*)
+  void Func2();
+  // expected-error@+3{{OpenACC clause 'bind' after 'device_type' clause on a 'routine' conflicts with the 'bind' clause before the first 'device_type'}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq bind(asdf) device_type(*) bind(asdf)
+  void Func3();
+  // expected-error@+3{{OpenACC clause 'bind' after 'device_type' clause on a 'routine' conflicts with the 'bind' clause before the first 'device_type'}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq bind(asdf) device_type(*) bind(asdf) device_type(*)
+  void Func4();
+  // expected-error@+3{{OpenACC clause 'bind' on a 'routine' directive conflicts with the 'bind' clause applying to the same 'device_type'}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq device_type(*) bind(asdf) bind(asdf)
+  void Func5();
+  // expected-error@+3{{OpenACC clause 'bind' on a 'routine' directive conflicts with the 'bind' clause applying to the same 'device_type'}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq device_type(*) bind(asdf) bind(asdf) device_type(*)
+  void Func6();
+  // expected-error@+3{{OpenACC clause 'bind' on a 'routine' directive conflicts with the 'bind' clause applying to the same 'device_type'}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc routine seq device_type(*) bind(asdf) device_type(*) bind(asdf) bind(asdf)
+  void Func7();
+
+#pragma acc routine seq device_type(*) bind(asdf) device_type(*) bind(asdf)
+  void Func8();
+}