[OpenACC] Fix bug with worker/vector/gang clause inside another

The original implementation rejected some valid constructs.  The rule is
supposed to be:

Gang-on-Kernel cannot have a gang in its region

Worker cannot have a worker or gang in its region

Vector cannot have worker, gang, or vector in its region.

The previous implementation improperly implemented that vector wasnt'
allowed in the other two.  This patch fixes it and adds testing for it.

Author: erichkeane

clang/lib/Sema/SemaOpenACC.cpp

@@ -1077,34 +1077,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
     }
   }
 
-  // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
-  // construct, the gang clause behaves as follows. ... The region of a loop
-  // with a gang clause may not contain another loop with a gang clause unless
-  // within a nested compute region.
-  if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) {
-    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
-    // one of these until we get to the end of the construct.
-    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
-        << OpenACCClauseKind::Vector << OpenACCClauseKind::Gang
-        << /*skip kernels construct info*/ 0;
-    SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
-  }
-
-  // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
-  // contain a loop with a gang or worker clause unless within a nested compute
-  // region.
-  if (SemaRef.LoopWorkerClauseLoc.isValid()) {
-    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
-    // one of these until we get to the end of the construct.
-    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
-        << OpenACCClauseKind::Vector << OpenACCClauseKind::Worker
-        << /*skip kernels construct info*/ 0;
-    SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
-  }
   // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
   // contain a loop with a gang, worker, or vector clause unless within a nested
   // compute region.

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

@@ -284,6 +284,13 @@ void Kernels() {
     for(int i = 0; i < 5; ++i);
   }
 
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+  for(int j = 0; j < 5; ++j) {
+#pragma acc loop worker(1) vector(1)
+    for(int i = 0; i < 5; ++i);
+  }
+
 #pragma acc kernels
 #pragma acc loop gang(num:1)
   for(int j = 0; j < 5; ++j) {

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

@@ -138,6 +138,12 @@ void uses() {
   for(int k = 0; k < 5; ++k);
   }
 
+#pragma acc loop worker
+  for(int i= 0; i< 5; ++i) {
+#pragma acc loop vector
+  for(int k = 0; k < 5; ++k);
+  }
+
 #pragma acc loop worker
   for(int i= 0; i< 5; ++i) {
 #pragma acc parallel