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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
