correctly diagnose duplicate 'tile' and 'collapse' clauses

After discussion, it is clear that we should do some funny-business
around 'device_type' on 'tile' and 'collapse', so this patch implements
that checking.

Author: erichkeane

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12921,6 +12921,9 @@ def err_acc_clause_routine_cannot_combine_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_clause_since_last_device_type
+    : Error<"OpenACC '%0' clause cannot appear more than once on a '%1' "
+            "directive %select{|in a 'device_type' region}2}">;
 
 def err_acc_reduction_num_gangs_conflict
     : Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not "

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -286,6 +286,33 @@ class SemaOpenACCClauseVisitor {
     return true;
   }
 
+  // For 'tile' and 'collapse', only allow 1 per 'device_type'.
+  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>);
+
+    if (Last == LastDeviceTypeItr)
+      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);
+
+    if (LastDeviceTypeItr != ExistingClauses.rend())
+      SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(),
+                   diag::note_acc_previous_clause_here);
+
+    // TODO: DIAG
+    return true;
+  }
+
 public:
   SemaOpenACCClauseVisitor(SemaOpenACC &S,
                            ArrayRef<const OpenACCClause *> ExistingClauses)
@@ -342,6 +369,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause(
 OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
 
+  if (DisallowSinceLastDeviceType<OpenACCTileClause>(Clause))
+    return nullptr;
+
   llvm::SmallVector<Expr *> NewSizeExprs;
 
   // Make sure these are all positive constant expressions or *.
@@ -1503,6 +1533,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
 OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
 
+  if (DisallowSinceLastDeviceType<OpenACCCollapseClause>(Clause))
+    return nullptr;
+
   ExprResult LoopCount = SemaRef.CheckCollapseLoopCount(Clause.getLoopCount());
 
   if (!LoopCount.isUsable())

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

@@ -14,6 +14,13 @@ void only_for_loops() {
 
 }
 
+void only_one_on_loop() {
+  // expected-error@+2{{OpenACC 'collapse' clause cannot appear more than once on a 'loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop collapse(1) collapse(1)
+  for(int i = 0; i < 5; ++i);
+}
+
 constexpr int three() { return 3; }
 constexpr int one() { return 1; }
 constexpr int neg() { return -1; }
@@ -497,12 +504,12 @@ template<unsigned N>
 void allow_multiple_collapse_templ() {
     // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
     // expected-note@+1{{active 'collapse' clause defined here}}
-#pragma acc loop collapse(N) collapse(N+1)
+#pragma acc loop collapse(N) device_type(*) collapse(N+1)
   for(unsigned i = 0; i < 5; ++i)
     for(unsigned j = 0; j < 5; ++j);
     // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
     // expected-note@+1{{active 'collapse' clause defined here}}
-#pragma acc loop collapse(N+1) collapse(N)
+#pragma acc loop collapse(N+1) device_type(*) collapse(N)
   for(unsigned i = 0; i < 5; ++i)
     for(unsigned j = 0; j < 5; ++j);
 }
@@ -511,12 +518,25 @@ void allow_multiple_collapse() {
 
     // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
     // expected-note@+1{{active 'collapse' clause defined here}}
-#pragma acc loop collapse(2) collapse(3)
+#pragma acc loop collapse(2) device_type(*) collapse(3)
   for(unsigned i = 0; i < 5; ++i)
     for(unsigned j = 0; j < 5; ++j);
     // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
     // expected-note@+1{{active 'collapse' clause defined here}}
-#pragma acc loop collapse(3) collapse(2)
+#pragma acc loop collapse(3) device_type(*) collapse(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+}
+
+void no_dupes_since_last_device_type() {
+  // expected-error@+3{{OpenACC 'collapse' clause cannot appear more than once on a 'loop' directive in a 'device_type' region}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop collapse(1) device_type(*) collapse(1) collapse(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+ 
+#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);
 }

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

@@ -143,6 +143,13 @@ void only_for_loops() {
     do{}while(true);
 }
 
+void only_one_on_loop() {
+  // expected-error@+2{{OpenACC 'tile' clause cannot appear more than once on a 'loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop tile(1) tile(1)
+  for(int i = 0; i < 5; ++i);
+}
+
 template<unsigned Val>
 void depth_too_high_templ() {
   // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
@@ -380,13 +387,13 @@ void intervening() {
 void use_largest_tile() {
 // expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}}
 // expected-note@+1{{active 'tile' clause defined here}}
-#pragma acc loop tile(1,2) tile (3,4,5)
+#pragma acc loop tile(1,2) device_type(*) tile (3,4,5)
   for(int i = 0; i < 5; ++i)
     for (int j = 0; j < 5; ++j);
 
 // expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}}
 // expected-note@+1{{active 'tile' clause defined here}}
-#pragma acc loop tile (3,4,5) tile(1,2)
+#pragma acc loop tile (3,4,5) device_type(*) tile(1,2)
   for(int i = 0; i < 5; ++i)
     for (int j = 0; j < 5; ++j);
 }
@@ -401,3 +408,15 @@ void collapse_tile_depth() {
     for(int j = 0; j < 5; ++j);
   }
 }
+void no_dupes_since_last_device_type() {
+  // expected-error@+3{{OpenACC 'tile' clause cannot appear more than once on a 'loop' directive in a 'device_type' region}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc loop tile(1) device_type(*) tile(1) tile(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+
+#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);
+}