Implement 'new' num_gangs/worker/vector_length rules around device_type, lots of test updates, fixed up the diagnostics-sema-kinds

Author: erichkeane

clang/include/clang/Basic/DiagnosticSemaKinds.td

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

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -334,6 +334,7 @@ class SemaOpenACCClauseVisitor {
   }
 
   // For 'tile' and 'collapse', only allow 1 per 'device_type'.
+  // Also applies to num_worker, num_gangs, and vector_length.
   template <typename TheClauseTy>
   bool DisallowSinceLastDeviceType(SemaOpenACC::OpenACCParsedClause &Clause) {
     auto LastDeviceTypeItr =
@@ -356,7 +357,6 @@ class SemaOpenACCClauseVisitor {
       SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(),
                    diag::note_acc_previous_clause_here);
 
-    // TODO: DIAG
     return true;
   }
 
@@ -483,6 +483,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause(
 OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
 
+  if (DisallowSinceLastDeviceType<OpenACCNumGangsClause>(Clause))
+    return nullptr;
+
   // num_gangs requires at least 1 int expr in all forms.  Diagnose here, but
   // allow us to continue, an empty clause might be useful for future
   // diagnostics.
@@ -573,6 +576,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
 OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
 
+  if (DisallowSinceLastDeviceType<OpenACCNumWorkersClause>(Clause))
+    return nullptr;
+
   // OpenACC 3.3 Section 2.9.2:
   // An argument is allowed only when the 'num_workers' does not appear on the
   // kernels construct.
@@ -601,6 +607,10 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
+
+  if (DisallowSinceLastDeviceType<OpenACCVectorLengthClause>(Clause))
+    return nullptr;
+
   // OpenACC 3.3 Section 2.9.4:
   // An argument is allowed only when the 'vector_length' does not appear on the
   // 'kernels' construct.

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

@@ -13,6 +13,31 @@ void Test() {
 #pragma acc parallel loop num_gangs(1)
   for(int i = 5; i < 10;++i);
 
+  // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(1) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop num_gangs(1) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'kernels loop' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(1) device_type(*) num_gangs(1) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'parallel loop' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop device_type(*) num_gangs(1) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
+#pragma acc parallel loop num_gangs(1) device_type(*) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
   // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type}}
 #pragma acc parallel loop num_gangs(getF())
   for(int i = 5; i < 10;++i);
@@ -25,12 +50,6 @@ void Test() {
 #pragma acc parallel loop num_gangs()
   for(int i = 5; i < 10;++i);
 
-#pragma acc kernels loop num_gangs(1) num_gangs(2)
-  for(int i = 5; i < 10;++i);
-
-#pragma acc parallel loop num_gangs(1) num_gangs(2)
-  for(int i = 5; i < 10;++i);
-
   // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels loop' directive expects maximum of 1, 2 were provided}}
 #pragma acc kernels loop num_gangs(1, getS())
   for(int i = 5; i < 10;++i);

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

@@ -13,6 +13,31 @@ void Test() {
 #pragma acc parallel loop num_workers(1)
   for(int i = 5; i < 10;++i);
 
+  // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'kernels loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop num_workers(1) num_workers(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'parallel loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop num_workers(1) num_workers(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'kernels loop' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop num_workers(1) device_type(*) num_workers(1) num_workers(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'parallel loop' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop device_type(*) num_workers(1) num_workers(2)
+  for(int i = 5; i < 10;++i);
+
+#pragma acc parallel loop num_workers(1) device_type(*) num_workers(2)
+  for(int i = 5; i < 10;++i);
+
   // expected-error@+1{{OpenACC clause 'num_workers' requires expression of integer type}}
 #pragma acc parallel loop num_workers(getF())
   for(int i = 5; i < 10;++i);

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

@@ -13,6 +13,31 @@ void Test() {
 #pragma acc parallel loop vector_length(1)
   for(int i = 5; i < 10;++i);
 
+  // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'kernels loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop vector_length(1) vector_length(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'parallel loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop vector_length(1) vector_length(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'kernels loop' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop vector_length(1) device_type(*) vector_length(1) vector_length(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'parallel loop' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop device_type(*) vector_length(1) vector_length(2)
+  for(int i = 5; i < 10;++i);
+
+#pragma acc parallel loop vector_length(1) device_type(*) vector_length(2)
+  for(int i = 5; i < 10;++i);
+
   // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type}}
 #pragma acc parallel loop vector_length(getF())
   for(int i = 5; i < 10;++i);

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

@@ -12,12 +12,31 @@ void Test() {
 #pragma acc parallel num_gangs(1)
   while(1);
 
+  // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels' directive}}
+  // expected-note@+1{{previous clause is here}}
 #pragma acc kernels num_gangs(1) num_gangs(2)
   while(1);
 
+  // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel' directive}}
+  // expected-note@+1{{previous clause is here}}
 #pragma acc parallel num_gangs(1) num_gangs(2)
   while(1);
 
+  // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_gangs(1) device_type(*) num_gangs(1) num_gangs(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(*) num_gangs(1) num_gangs(2)
+  while(1);
+
+#pragma acc parallel num_gangs(1) device_type(*) num_gangs(2)
+  while(1);
+
   // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels' directive expects maximum of 1, 2 were provided}}
 #pragma acc kernels num_gangs(1, getS())
   while(1);

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

@@ -8,6 +8,31 @@ void Test() {
 #pragma acc kernels num_workers(1)
   while(1);
 
+  // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'kernels' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_workers(1) num_workers(2)
+  while(1);
+
+  // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'parallel' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel num_workers(1) num_workers(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_workers(1) device_type(*) num_workers(1) num_workers(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(*) num_workers(1) num_workers(2)
+  while(1);
+
+#pragma acc parallel num_workers(1) device_type(*) num_workers(2)
+  while(1);
+
   // expected-error@+1{{OpenACC 'num_workers' clause is not valid on 'serial' directive}}
 #pragma acc serial num_workers(1)
   while(1);

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

@@ -8,6 +8,31 @@ void Test() {
 #pragma acc kernels vector_length(1)
   while(1);
 
+  // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'kernels' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels vector_length(1) vector_length(2)
+  while(1);
+
+  // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'parallel' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel vector_length(1) vector_length(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels vector_length(1) device_type(*) vector_length(1) vector_length(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel device_type(*) vector_length(1) vector_length(2)
+  while(1);
+
+#pragma acc parallel vector_length(1) device_type(*) vector_length(2)
+  while(1);
+
   // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial' directive}}
 #pragma acc serial vector_length(1)
   while(1);

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

@@ -529,7 +529,7 @@ void allow_multiple_collapse() {
 }
 
 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-error@+3{{OpenACC 'collapse' clause cannot appear more than once in a 'device_type' region on a 'loop' directive}}
   // 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)

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

@@ -409,7 +409,7 @@ void collapse_tile_depth() {
   }
 }
 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-error@+3{{OpenACC 'tile' clause cannot appear more than once in a 'device_type' region on a 'loop' directive}}
   // 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)