Correct implementation of 'device_type' on routine

After looking into it further, the rule for the 'worker', 'gang',
'vector', and 'seq' on a 'routine' construct are suppsoed to be more
complicated.

The area BEFORE the first device_type are 'global' and apply to all
following device_types.

Then, each device_type region can name their own.

The rule is it is exclusive on a per-device_type basis for these.

Author: erichkeane

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12912,6 +12912,16 @@ def err_acc_clause_after_device_type
 def err_acc_clause_cannot_combine
     : Error<"OpenACC clause '%0' may not appear on the same construct as a "
             "'%1' clause on a '%2' construct">;
+def err_acc_clause_routine_cannot_combine_before_device_type
+    : Error<"OpenACC clause '%0' after 'device_type' clause on a 'routine' "
+              "conflicts with the '%1' clause before the first 'device_type'">;
+def err_acc_clause_routine_cannot_combine_same_device_type
+    : Error<"OpenACC clause '%0' on a 'routine' directive conflicts with the "
+            "'%1' clause applying to the same 'device_type'">;
+def err_acc_clause_routine_one_of_in_region
+    : Error<"OpenACC 'routine' construct must have at least one 'gang', 'seq', "
+            "'vector', or 'worker' clause that applies to each 'device_type'">;
+
 def err_acc_reduction_num_gangs_conflict
     : Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not "
             "appear on a '%2' construct "
@@ -12980,9 +12990,10 @@ def err_acc_gang_reduction_numgangs_conflict
             "'%1' clause %select{inside a compute construct with a|and a}3 "
             "'num_gangs' clause with more than one argument">;
 
-    def err_reduction_op_mismatch
+def err_reduction_op_mismatch
     : Error<"OpenACC 'reduction' variable must have the same operator in all "
             "nested constructs (%0 vs %1)">;
+
 def err_acc_loop_variable_type
     : Error<"loop variable of loop associated with an OpenACC '%0' construct "
             "must be of integer, pointer, or random-access-iterator type (is "

clang/lib/Sema/SemaOpenACC.cpp

@@ -1828,6 +1828,62 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(
   llvm_unreachable("Invalid associated statement application");
 }
 
+namespace {
+
+// Routine has some pretty complicated set of rules for how device_type
+// interacts with 'gang', 'worker', 'vector', and 'seq'. Enforce  part of it
+// here.
+bool CheckValidRoutineGangWorkerVectorSeqClauses(
+    SemaOpenACC &SemaRef, SourceLocation DirectiveLoc,
+    ArrayRef<const OpenACCClause *> Clauses) {
+  auto RequiredPred = llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
+                                    OpenACCVectorClause, OpenACCSeqClause>;
+  // The clause handling has assured us that there is no duplicates.  That is,
+  // if there is 1 before a device_type, there are none after a device_type.
+  // If not, there is at most 1 applying to each device_type.
+
+  // What is left to legalize is that either:
+  // 1- there is 1 before the first device_type.
+  // 2- there is 1 AFTER each device_type.
+  auto *FirstDeviceType =
+      llvm::find_if(Clauses, llvm::IsaPred<OpenACCDeviceTypeClause>);
+
+  // If there is 1 before the first device_type (or at all if no device_type),
+  // we are legal.
+  auto *ClauseItr =
+      std::find_if(Clauses.begin(), FirstDeviceType, RequiredPred);
+
+  if (ClauseItr != FirstDeviceType)
+    return false;
+
+  // If there IS no device_type, and no clause, diagnose.
+  if (FirstDeviceType == Clauses.end())
+    return SemaRef.Diag(DirectiveLoc, diag::err_acc_construct_one_clause_of)
+           << OpenACCDirectiveKind::Routine
+           << "'gang', 'seq', 'vector', or 'worker'";
+
+  // Else, we have to check EACH device_type group. PrevDeviceType is the
+  // device-type before the current group.
+  auto *PrevDeviceType = FirstDeviceType;
+
+  while (PrevDeviceType != Clauses.end()) {
+    auto *NextDeviceType =
+        std::find_if(std::next(PrevDeviceType), Clauses.end(),
+                     llvm::IsaPred<OpenACCDeviceTypeClause>);
+
+    ClauseItr = std::find_if(PrevDeviceType, NextDeviceType, RequiredPred);
+
+    if (ClauseItr == NextDeviceType)
+      return SemaRef.Diag((*PrevDeviceType)->getBeginLoc(),
+                          diag::err_acc_clause_routine_one_of_in_region);
+
+    PrevDeviceType = NextDeviceType;
+  }
+
+  return false;
+}
+} // namespace
+
 bool SemaOpenACC::ActOnStartDeclDirective(
     OpenACCDirectiveKind K, SourceLocation StartLoc,
     ArrayRef<const OpenACCClause *> Clauses) {
@@ -1839,6 +1895,9 @@ bool SemaOpenACC::ActOnStartDeclDirective(
 
   if (DiagnoseRequiredClauses(K, StartLoc, Clauses))
     return true;
+  if (K == OpenACCDirectiveKind::Routine &&
+      CheckValidRoutineGangWorkerVectorSeqClauses(*this, StartLoc, Clauses))
+    return true;
 
   return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/false);
 }

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -199,6 +199,93 @@ 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) {
+    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);
+
+      if (ExistingProhibitedClause == ExistingClauses.end())
+        return false;
+
+      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
+          << Clause.getClauseKind()
+          << (*ExistingProhibitedClause)->getClauseKind()
+          << Clause.getDirectiveKind();
+      SemaRef.Diag((*ExistingProhibitedClause)->getBeginLoc(),
+                   diag::note_acc_previous_clause_here);
+      return true;
+    }
+
+    // At this point we know that this is 'after' a device type. So this is an
+    // error if: 1- there is one BEFORE the 'device_type' 2- there is one
+    // between this and the previous 'device_type'.
+
+    auto *BeforeDeviceType =
+        std::find_if(ExistingClauses.begin(), FirstDeviceType, ProhibitedPred);
+
+    // 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) {
+      SemaRef.Diag(
+          Clause.getBeginLoc(),
+          diag::err_acc_clause_routine_cannot_combine_before_device_type)
+          << Clause.getClauseKind() << (*BeforeDeviceType)->getClauseKind();
+      SemaRef.Diag((*BeforeDeviceType)->getBeginLoc(),
+                   diag::note_acc_previous_clause_here);
+      SemaRef.Diag((*FirstDeviceType)->getBeginLoc(),
+                   diag::note_acc_previous_clause_here);
+      return true;
+    }
+
+    auto LastDeviceTypeItr =
+        std::find_if(ExistingClauses.rbegin(), ExistingClauses.rend(),
+                     llvm::IsaPred<OpenACCDeviceTypeClause>);
+
+    // We already know there is one in the list, so it is nonsensical to not
+    // have one.
+    assert(LastDeviceTypeItr != ExistingClauses.rend());
+
+    // Get the device-type from-the-front (not reverse) iterator from the
+    // reverse iterator.
+    auto *LastDeviceType = LastDeviceTypeItr.base() - 1;
+
+    auto *ExistingProhibitedSinceLastDevice =
+        std::find_if(LastDeviceType, ExistingClauses.end(), ProhibitedPred);
+
+    // No prohibited ones since the last device-type.
+    if (ExistingProhibitedSinceLastDevice == ExistingClauses.end())
+      return false;
+
+    SemaRef.Diag(Clause.getBeginLoc(),
+                 diag::err_acc_clause_routine_cannot_combine_same_device_type)
+        << Clause.getClauseKind()
+        << (*ExistingProhibitedSinceLastDevice)->getClauseKind();
+    SemaRef.Diag((*ExistingProhibitedSinceLastDevice)->getBeginLoc(),
+                 diag::note_acc_previous_clause_here);
+    SemaRef.Diag((*LastDeviceType)->getBeginLoc(),
+                 diag::note_acc_previous_clause_here);
+    return true;
+  }
+
 public:
   SemaOpenACCClauseVisitor(SemaOpenACC &S,
                            ArrayRef<const OpenACCClause *> ExistingClauses)
@@ -213,6 +300,8 @@ class SemaOpenACCClauseVisitor {
                                          Clause.getClauseKind(),
                                          Clause.getBeginLoc(), ExistingClauses))
       return nullptr;
+    if (CheckValidRoutineGangWorkerVectorSeqNewClause(Clause))
+      return nullptr;
 
     switch (Clause.getClauseKind()) {
     case OpenACCClauseKind::Shortloop:

clang/test/ParserOpenACC/parse-constructs.c

@@ -153,20 +153,20 @@ void func() {
 #pragma acc routine seq
 void routine_func();
 // expected-error@+2{{invalid OpenACC clause 'clause'}}
-// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}}
+// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}}
 #pragma acc routine clause list
 void routine_func();
 
 // expected-error@+1{{use of undeclared identifier 'func_name'}}
 #pragma acc routine (func_name) seq
 // expected-error@+3{{use of undeclared identifier 'func_name'}}
 // expected-error@+2{{invalid OpenACC clause 'clause'}}
-// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}}
+// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}}
 #pragma acc routine (func_name) clause list
 
 #pragma acc routine (routine_func) seq
 // expected-error@+2{{invalid OpenACC clause 'clause'}}
-// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}}
+// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}}
 #pragma acc routine (routine_func) clause list
 
 // expected-error@+2{{expected ')'}}