[clang] cb6a02a - [OpenACC] Implement 'worker' clause for combined constructs
via cfe-commits
cfe-commits at lists.llvm.org
Thu Dec 5 10:43:56 PST 2024
Author: erichkeane
Date: 2024-12-05T10:43:52-08:00
New Revision: cb6a02abe21fb399e86863dd69e865d0ddaa6838
URL: https://github.com/llvm/llvm-project/commit/cb6a02abe21fb399e86863dd69e865d0ddaa6838
DIFF: https://github.com/llvm/llvm-project/commit/cb6a02abe21fb399e86863dd69e865d0ddaa6838.diff
LOG: [OpenACC] Implement 'worker' clause for combined constructs
This is very similar to 'gang', except with fewer restrictions, and only an
interaction with 'num_workers', plus disallowing 'gang' and 'worker' in
its associated statement. This patch implements this, the same as how
'gang' implemented it.
Added:
clang/test/SemaOpenACC/combined-construct-worker-ast.cpp
clang/test/SemaOpenACC/combined-construct-worker-clause.cpp
Modified:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaOpenACC.cpp
clang/test/AST/ast-print-openacc-combined-construct.cpp
clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
clang/test/SemaOpenACC/combined-construct-device_type-clause.c
clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 447358f0a5f382..a8fd11a56edc83 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12766,8 +12766,8 @@ def err_acc_num_arg_conflict
"construct%select{| associated with a '%3' construct}2 that has a "
"'%4' clause">;
def err_acc_num_arg_conflict_reverse
- : Error<"'num_gangs' clause not allowed on a 'kernels loop' construct that "
- "has a 'gang' clause with a 'num' argument">;
+ : Error<"'%0' clause not allowed on a 'kernels loop' construct that "
+ "has a '%1' clause with a%select{n| 'num'}2 argument">;
def err_acc_clause_in_clause_region
: Error<"loop with a '%0' clause may not exist in the region of a '%1' "
"clause%select{| on a '%3' construct}2">;
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 16348d0b8837ef..5fc45974fa0ddc 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -748,7 +748,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
for (auto *GC : GangClauses) {
if (cast<OpenACCGangClause>(GC)->hasExprOfKind(OpenACCGangKind::Num)) {
SemaRef.Diag(Clause.getBeginLoc(),
- diag::err_acc_num_arg_conflict_reverse);
+ diag::err_acc_num_arg_conflict_reverse)
+ << OpenACCClauseKind::NumGangs << OpenACCClauseKind::Gang
+ << /*Num argument*/ 1;
SemaRef.Diag(GC->getBeginLoc(), diag::note_acc_previous_clause_here);
return nullptr;
}
@@ -768,6 +770,25 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, 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.
+ if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) {
+ auto WorkerClauses = llvm::make_filter_range(
+ ExistingClauses, llvm::IsaPred<OpenACCWorkerClause>);
+
+ for (auto *WC : WorkerClauses) {
+ if (cast<OpenACCWorkerClause>(WC)->hasIntExpr()) {
+ SemaRef.Diag(Clause.getBeginLoc(),
+ diag::err_acc_num_arg_conflict_reverse)
+ << OpenACCClauseKind::NumWorkers << OpenACCClauseKind::Worker
+ << /*num argument*/ 0;
+ SemaRef.Diag(WC->getBeginLoc(), diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+ }
+ }
+
assert(Clause.getIntExprs().size() == 1 &&
"Invalid number of expressions for NumWorkers");
return OpenACCNumWorkersClause::Create(
@@ -1254,74 +1275,107 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
if (DiagIfSeqClause(Clause))
return nullptr;
- // Restrictions only properly implemented on 'loop' constructs, and it is
- // the only construct that can do anything with this, so skip/treat as
- // unimplemented for the combined constructs.
- if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
+ // Restrictions only properly implemented on 'loop'/'combined' constructs, and
+ // it is the only construct that can do anything with this, so skip/treat as
+ // unimplemented for the routine constructs.
+ if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop &&
+ !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
return isNotImplemented();
Expr *IntExpr =
Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
if (IntExpr) {
- switch (SemaRef.getActiveComputeConstructInfo().Kind) {
- case OpenACCDirectiveKind::Invalid:
- case OpenACCDirectiveKind::Parallel:
- case OpenACCDirectiveKind::Serial:
- DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
- OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
- SemaRef.getActiveComputeConstructInfo().Kind);
- IntExpr = nullptr;
- break;
- case OpenACCDirectiveKind::Kernels: {
- const auto *Itr =
- llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
- llvm::IsaPred<OpenACCNumWorkersClause>);
- if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
- << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
- << HasAssocKind(Clause.getDirectiveKind(),
- SemaRef.getActiveComputeConstructInfo().Kind)
- << SemaRef.getActiveComputeConstructInfo().Kind
- << OpenACCClauseKind::NumWorkers;
- SemaRef.Diag((*Itr)->getBeginLoc(),
- diag::note_acc_previous_clause_here);
-
+ if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ switch (SemaRef.getActiveComputeConstructInfo().Kind) {
+ case OpenACCDirectiveKind::Invalid:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::Serial:
+ DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+ OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind);
IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::KernelsLoop:
+ case OpenACCDirectiveKind::Kernels: {
+ const auto *Itr =
+ llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+ llvm::IsaPred<OpenACCNumWorkersClause>);
+ if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+ << HasAssocKind(Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind)
+ << SemaRef.getActiveComputeConstructInfo().Kind
+ << OpenACCClauseKind::NumWorkers;
+ SemaRef.Diag((*Itr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+
+ IntExpr = nullptr;
+ }
+ break;
+ }
+ default:
+ llvm_unreachable("Non compute construct in active compute construct");
+ }
+ } else {
+ if (Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop ||
+ Clause.getDirectiveKind() == OpenACCDirectiveKind::SerialLoop) {
+ DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+ OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind);
+ IntExpr = nullptr;
+ } else {
+ assert(Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop &&
+ "Unknown combined directive kind?");
+ const auto *Itr = llvm::find_if(ExistingClauses,
+ llvm::IsaPred<OpenACCNumWorkersClause>);
+ if (Itr != ExistingClauses.end()) {
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+ << HasAssocKind(Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind)
+ << SemaRef.getActiveComputeConstructInfo().Kind
+ << OpenACCClauseKind::NumWorkers;
+ SemaRef.Diag((*Itr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+
+ IntExpr = nullptr;
+ }
}
- break;
- }
- default:
- llvm_unreachable("Non compute construct in active compute construct");
}
}
- // 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::Worker << OpenACCClauseKind::Worker
- << /*skip kernels construct info*/ 0;
- SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
- diag::note_acc_previous_clause_here);
- return nullptr;
- }
+ if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ // 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::Worker << 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.
- if (SemaRef.LoopVectorClauseLoc.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::Worker << OpenACCClauseKind::Vector
- << /*skip kernels construct info*/ 0;
- SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
- 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.
+ if (SemaRef.LoopVectorClauseLoc.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::Worker << OpenACCClauseKind::Vector
+ << /*skip kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
}
return OpenACCWorkerClause::Create(Ctx, Clause.getBeginLoc(),
diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index 40f174e539c01e..45aaeb41a5ce0e 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -313,4 +313,35 @@ void foo() {
// CHECK-NEXT: ;
#pragma acc serial loop gang(static:*)
for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop worker
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop worker
+ for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc parallel loop worker
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop worker
+ for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc serial loop worker
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop worker
+ for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop worker(num: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop worker(5)
+ for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop worker(num: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop worker(num:5)
+ for(int i = 0;i<5;++i);
+
}
diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
index e35bd6da2f18b3..a675db2a70a12f 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -46,7 +46,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc parallel loop auto if_present
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc parallel loop auto worker
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -166,7 +165,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc parallel loop if_present auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc parallel loop worker auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -287,7 +285,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc parallel loop independent if_present
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc parallel loop independent worker
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -407,7 +404,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc parallel loop if_present independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc parallel loop worker independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -650,9 +646,8 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc parallel loop gang seq
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}}
- // TODOexpected-note at +1{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc parallel loop worker seq
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}}
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index d3ed8234d16b14..3bf88d09c65e64 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -57,7 +57,6 @@ void uses() {
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop device_type(*) auto
for(int i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc parallel loop device_type(*) worker
for(int i = 0; i < 5; ++i);
// expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'serial loop' construct}}
diff --git a/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
index f179b928215e71..167f56289c58dd 100644
--- a/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
+++ b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
@@ -1,4 +1,3 @@
-
// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
// Test this with PCH.
diff --git a/clang/test/SemaOpenACC/combined-construct-worker-ast.cpp b/clang/test/SemaOpenACC/combined-construct-worker-ast.cpp
new file mode 100644
index 00000000000000..6ac0fdb35709fb
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-worker-ast.cpp
@@ -0,0 +1,135 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+#ifndef PCH_HELPER
+#define PCH_HELPER
+void NormalUses() {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK-NEXT: CompoundStmt
+
+ int Val;
+
+#pragma acc parallel loop worker
+ for(int i = 0; i < 5; ++i);
+ // CHECK: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+#pragma acc serial loop worker
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc kernels loop worker(Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+#pragma acc kernels loop worker(num:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+}
+
+template<typename T, unsigned One>
+void TemplateUses(T Val) {
+ // CHECK: FunctionTemplateDecl{{.*}}TemplateUses
+ // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 One
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplateUses 'void (T)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} referenced Val 'T'
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop worker
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+#pragma acc serial loop worker
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc kernels loop worker(Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+#pragma acc kernels loop worker(num:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc kernels loop worker(num:One)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplateUses 'void (int)' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: TemplateArgument integral '1U'
+ // CHECK-NEXT: ParmVarDecl{{.*}} used Val 'int'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: worker clause
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+ // CHECK-NEXT: NonTypeTemplateParmDecl
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+}
+
+void inst() {
+ TemplateUses<int, 1>(5);
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-worker-clause.cpp b/clang/test/SemaOpenACC/combined-construct-worker-clause.cpp
new file mode 100644
index 00000000000000..1351f9ea0c69c4
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-worker-clause.cpp
@@ -0,0 +1,150 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<unsigned I>
+void TemplUses() {
+
+#pragma acc parallel loop worker
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc parallel loop worker()
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop worker(I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc serial loop worker
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc serial loop worker()
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'serial loop' construct}}
+#pragma acc serial loop worker(I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'serial loop' construct}}
+#pragma acc serial loop worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop worker(I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num' argument to 'worker' clause not allowed on a 'kernels loop' construct that has a 'num_workers' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_workers(1) worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num_workers' clause not allowed on a 'kernels loop' construct that has a 'worker' clause with an argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop worker(num:I) num_workers(1)
+ for(int i = 0; i < 5; ++i);
+}
+
+void NormalUses() {
+ TemplUses<4>();
+
+ int I;
+#pragma acc parallel loop worker
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc parallel loop worker()
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop worker(I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc serial loop worker
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc serial loop worker()
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'serial loop' construct}}
+#pragma acc serial loop worker(I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'serial loop' construct}}
+#pragma acc serial loop worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop worker(I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num' argument to 'worker' clause not allowed on a 'kernels loop' construct that has a 'num_workers' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_workers(1) worker(num:I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num_workers' clause not allowed on a 'kernels loop' construct that has a 'worker' clause with an argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop worker(num:I) num_workers(1)
+ for(int i = 0; i < 5; ++i);
+
+ // OK, kernels loop is a new compute construct
+#pragma acc kernels num_workers(1)
+ for(int i = 0; i < 5; ++i) {
+#pragma acc kernels loop worker(num:1)
+ for(int i = 0; i < 5; ++i);
+ }
+#pragma acc kernels loop num_workers(1)
+ for(int i = 0; i < 5; ++i) {
+#pragma acc kernels loop worker(num:1)
+ for(int i = 0; i < 5; ++i);
+ }
+
+#pragma acc kernels loop num_workers(1)
+ for(int i = 0; i < 5; ++i) {
+ // expected-error at +2{{'num' argument to 'worker' clause not allowed on a 'loop' construct associated with a 'kernels loop' construct that has a 'num_workers' clause}}
+ // expected-note at -3{{previous clause is here}}
+#pragma acc loop worker(num:1)
+ for(int i = 0; i < 5; ++i);
+ }
+
+#pragma acc parallel loop worker
+ for(int i = 0; i < 5; ++i) {
+ // expected-error at +4{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+ // expected-note at -3{{previous clause is here}}
+ // expected-error at +2{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+ // expected-note at -5{{previous clause is here}}
+#pragma acc loop gang, worker, vector
+ for(int i = 0; i < 5; ++i);
+ }
+#pragma acc kernels loop worker
+ for(int i = 0; i < 5; ++i) {
+ // expected-error at +4{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+ // expected-note at -3{{previous clause is here}}
+ // expected-error at +2{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+ // expected-note at -5{{previous clause is here}}
+#pragma acc loop gang, worker, vector
+ for(int i = 0; i < 5; ++i);
+ }
+#pragma acc serial loop worker
+ for(int i = 0; i < 5; ++i) {
+ // expected-error at +4{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+ // expected-note at -3{{previous clause is here}}
+ // expected-error at +2{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+ // expected-note at -5{{previous clause is here}}
+#pragma acc loop gang, worker, vector
+ for(int i = 0; i < 5; ++i);
+ }
+}
More information about the cfe-commits
mailing list