[OpenACC] Implement better dupe catching for device_type clauses

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13075,6 +13075,9 @@ def err_acc_clause_routine_one_of_in_region
 def err_acc_clause_since_last_device_type
     : Error<"OpenACC '%0' clause cannot appear more than once%select{| in a "
             "'device_type' region}2 on a '%1' directive">;
+def err_acc_clause_conflicts_prev_dev_type
+    : Error<"OpenACC '%0' clause applies to 'device_type' '%1', which "
+            "conflicts with previous clause">;
 
 def err_acc_reduction_num_gangs_conflict
     : Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not "

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -335,29 +335,103 @@ class SemaOpenACCClauseVisitor {
 
   // For 'tile' and 'collapse', only allow 1 per 'device_type'.
   // Also applies to num_worker, num_gangs, vector_length, and async.
+  // This does introspection into the actual device-types to prevent duplicates
+  // across device types as well.
   template <typename TheClauseTy>
   bool DisallowSinceLastDeviceType(SemaOpenACC::OpenACCParsedClause &Clause) {
     auto LastDeviceTypeItr =
         std::find_if(ExistingClauses.rbegin(), ExistingClauses.rend(),
                      llvm::IsaPred<OpenACCDeviceTypeClause>);
 
-    auto Last = std::find_if(ExistingClauses.rbegin(), LastDeviceTypeItr,
-                             llvm::IsaPred<TheClauseTy>);
+    auto LastSinceDevTy =
+        std::find_if(ExistingClauses.rbegin(), LastDeviceTypeItr,
+                     llvm::IsaPred<TheClauseTy>);
 
-    if (Last == LastDeviceTypeItr)
+    // In this case there is a duplicate since the last device_type/lack of a
+    // device_type.  Diagnose these as duplicates.
+    if (LastSinceDevTy != LastDeviceTypeItr) {
+      SemaRef.Diag(Clause.getBeginLoc(),
+                   diag::err_acc_clause_since_last_device_type)
+          << Clause.getClauseKind() << Clause.getDirectiveKind()
+          << (LastDeviceTypeItr != ExistingClauses.rend());
+      SemaRef.Diag((*LastSinceDevTy)->getBeginLoc(),
+                   diag::note_acc_previous_clause_here);
+
+      // Mention the last device_type as well.
+      if (LastDeviceTypeItr != ExistingClauses.rend())
+        SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(),
+                     diag::note_acc_previous_clause_here);
+      return true;
+    }
+
+    // If this isn't in a device_type, and we didn't diagnose that there are
+    // dupes above, just give up, no sense in searching for previous device_type
+    // regions as they don't exist.
+    if (LastDeviceTypeItr == ExistingClauses.rend())
       return false;
 
-    SemaRef.Diag(Clause.getBeginLoc(),
-                 diag::err_acc_clause_since_last_device_type)
-        << Clause.getClauseKind() << Clause.getDirectiveKind()
-        << (LastDeviceTypeItr != ExistingClauses.rend());
-    SemaRef.Diag((*Last)->getBeginLoc(), diag::note_acc_previous_clause_here);
+    // The device-type that is active for us, so we can compare to the previous
+    // ones.
+    const auto &ActiveDeviceTypeClause =
+        cast<OpenACCDeviceTypeClause>(**LastDeviceTypeItr);
 
-    if (LastDeviceTypeItr != ExistingClauses.rend())
-      SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(),
-                   diag::note_acc_previous_clause_here);
+    auto PrevDeviceTypeItr = LastDeviceTypeItr;
+    auto CurDevTypeItr = LastDeviceTypeItr;
 
-    return true;
+    while ((CurDevTypeItr = std::find_if(
+                std::next(PrevDeviceTypeItr), ExistingClauses.rend(),
+                llvm::IsaPred<OpenACCDeviceTypeClause>)) !=
+           ExistingClauses.rend()) {
+      // At this point, we know that we have a region between two device_types,
+      // as specified by CurDevTypeItr and PrevDeviceTypeItr.
+
+      auto CurClauseKindItr = std::find_if(PrevDeviceTypeItr, CurDevTypeItr,
+                                           llvm::IsaPred<TheClauseTy>);
+
+      // There are no clauses of the current kind between these device_types, so
+      // continue.
+      if (CurClauseKindItr == CurDevTypeItr)
+        continue;
+
+      // At this point, we know that this device_type region has a collapse.  So
+      // diagnose if the two device_types have any overlap in their
+      // architectures.
+      const auto &CurDeviceTypeClause =
+          cast<OpenACCDeviceTypeClause>(**CurDevTypeItr);
+
+      for (const DeviceTypeArgument &arg :
+           ActiveDeviceTypeClause.getArchitectures()) {
+        for (const DeviceTypeArgument &prevArg :
+             CurDeviceTypeClause.getArchitectures()) {
+
+          // This should catch duplicates * regions, duplicate same-text (thanks
+          // to identifier equiv.) and case insensitive dupes.
+          if (arg.getIdentifierInfo() == prevArg.getIdentifierInfo() ||
+              (arg.getIdentifierInfo() && prevArg.getIdentifierInfo() &&
+               StringRef{arg.getIdentifierInfo()->getName()}.equals_insensitive(
+                   prevArg.getIdentifierInfo()->getName()))) {
+            SemaRef.Diag(Clause.getBeginLoc(),
+                         diag::err_acc_clause_conflicts_prev_dev_type)
+                << Clause.getClauseKind()
+                << (arg.getIdentifierInfo() ? arg.getIdentifierInfo()->getName()
+                                            : "*");
+            // mention the active device type.
+            SemaRef.Diag(ActiveDeviceTypeClause.getBeginLoc(),
+                         diag::note_acc_previous_clause_here);
+            // mention the previous clause.
+            SemaRef.Diag((*CurClauseKindItr)->getBeginLoc(),
+                         diag::note_acc_previous_clause_here);
+            // mention the previous device type.
+            SemaRef.Diag(CurDeviceTypeClause.getBeginLoc(),
+                         diag::note_acc_previous_clause_here);
+            return true;
+          }
+        }
+      }
+
+      PrevDeviceTypeItr = CurDevTypeItr;
+    }
+    return false;
   }
 
 public:

clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c

@@ -71,3 +71,18 @@ void Test() {
 #pragma acc loop num_gangs(1)
   for(int i = 5; i < 10;++i);
 }
+
+void no_dupes_since_last_device_type() {
+  // expected-error@+4{{OpenACC 'num_gangs' clause applies to 'device_type' 'radeon', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia, radeon) num_gangs(getS()) device_type(radeon) num_gangs(getS())
+  ;
+  // expected-error@+4{{OpenACC 'num_gangs' clause applies to 'device_type' 'nvidia', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia) num_gangs(getS()) device_type(nvidia, radeon) num_gangs(getS())
+  ;
+}

clang/test/SemaOpenACC/compute-construct-num_workers-clause.c

@@ -60,3 +60,18 @@ void Test() {
 #pragma acc loop num_workers(1)
   for(int i = 5; i < 10;++i);
 }
+
+void no_dupes_since_last_device_type() {
+  // expected-error@+4{{OpenACC 'num_workers' clause applies to 'device_type' 'radEon', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia, radeon) num_workers(getS()) device_type(radEon) num_workers(getS())
+  ;
+  // expected-error@+4{{OpenACC 'num_workers' clause applies to 'device_type' 'nvidia', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia) num_workers(getS()) device_type(nvidia, radeon) num_workers(getS())
+  ;
+}

clang/test/SemaOpenACC/compute-construct-vector_length-clause.c

@@ -60,3 +60,18 @@ void Test() {
 #pragma acc loop vector_length(1)
   for(int i = 5; i < 10;++i);
 }
+
+void no_dupes_since_last_device_type() {
+  // expected-error@+4{{OpenACC 'vector_length' clause applies to 'device_type' 'radeon', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia, radeon) vector_length(getS()) device_type(radeon) vector_length(getS())
+  ;
+  // expected-error@+4{{OpenACC 'vector_length' clause applies to 'device_type' 'nvidia', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(nvidia) vector_length(getS()) device_type(nvidia, radeon) vector_length(getS())
+  ;
+}

clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp

@@ -539,4 +539,25 @@ void no_dupes_since_last_device_type() {
 #pragma acc loop collapse(1) device_type(*) collapse(1) device_type(nvidia) collapse(2)
   for(unsigned i = 0; i < 5; ++i)
     for(unsigned j = 0; j < 5; ++j);
+
+  // This one is ok, despite * being the 'all' value.
+#pragma acc loop device_type(*) collapse(1) device_type(nvidia) collapse(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+
+  // expected-error@+4{{OpenACC 'collapse' clause applies to 'device_type' 'nvidia', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(nviDia, radeon) collapse(1) device_type(nvidia) collapse(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+
+  // expected-error@+4{{OpenACC 'collapse' clause applies to 'device_type' 'radeon', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(radeon) collapse(1) device_type(nvidia, radeon) collapse(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
 }

clang/test/SemaOpenACC/loop-construct-tile-clause.cpp

@@ -419,4 +419,20 @@ void no_dupes_since_last_device_type() {
 #pragma acc loop tile(1) device_type(*) tile(1) device_type(nvidia) tile(2)
   for(unsigned i = 0; i < 5; ++i)
     for(unsigned j = 0; j < 5; ++j);
+
+  // expected-error@+4{{OpenACC 'tile' clause applies to 'device_type' 'nvidiA', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(nvidia, radeon) tile(1) device_type(nvidiA) tile(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+
+  // expected-error@+4{{OpenACC 'tile' clause applies to 'device_type' 'radeon', which conflicts with previous clause}}
+  // expected-note@+3{{previous clause is here}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop device_type(radeon) tile(1) device_type(nvidia, radeon) tile(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
 }