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

[via cfe-commits]

Author: erichkeane
Date: 2024-12-04T10:34:50-08:00
New Revision: e3446b9a079f1d911c96ae84d899d9ccd55a9951

URL: https://github.com/llvm/llvm-project/commit/e3446b9a079f1d911c96ae84d899d9ccd55a9951
DIFF: https://github.com/llvm/llvm-project/commit/e3446b9a079f1d911c96ae84d899d9ccd55a9951.diff

LOG: [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.

Added: 
    

Modified: 
    clang/lib/Sema/SemaOpenACC.cpp
    clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
    clang/test/SemaOpenACC/loop-construct-worker-clause.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 060df967322ac5..30491d6f3a3bd6 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/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.

diff  --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
index a22e7e865a801f..a2bda5a7e82f63 100644
--- a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
+++ b/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) {

diff  --git a/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp b/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
index 8b95981c2b80fe..35c63e467a5058 100644
--- a/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
+++ b/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