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@+1{{OpenACC clause 'if_present' not yet implemented}} #pragma acc parallel loop auto if_present for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}} #pragma acc parallel loop auto worker for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} @@ -166,7 +165,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}} #pragma acc parallel loop if_present auto for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}} #pragma acc parallel loop worker auto for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} @@ -287,7 +285,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}} #pragma acc parallel loop independent if_present for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}} #pragma acc parallel loop independent worker for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} @@ -407,7 +404,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}} #pragma acc parallel loop if_present independent for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}} #pragma acc parallel loop worker independent for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} @@ -650,9 +646,8 @@ void uses() { // expected-note@+1{{previous clause is here}} #pragma acc parallel loop gang seq for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}} - // TODOexpected-note@+1{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}} + // expected-note@+1{{previous clause is here}} #pragma acc parallel loop worker seq for(unsigned i = 0; i < 5; ++i); // TODOexpected-error@+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@+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@+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@+1{{expected expression}} +#pragma acc parallel loop worker() + for(int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+1{{expected expression}} +#pragma acc serial loop worker() + for(int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+2{{'num' argument to 'worker' clause not allowed on a 'kernels loop' construct that has a 'num_workers' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_workers(1) worker(num:I) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num_workers' clause not allowed on a 'kernels loop' construct that has a 'worker' clause with an argument}} + // expected-note@+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@+1{{expected expression}} +#pragma acc parallel loop worker() + for(int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+1{{expected expression}} +#pragma acc serial loop worker() + for(int i = 0; i < 5; ++i); + + // expected-error@+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@+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@+2{{'num' argument to 'worker' clause not allowed on a 'kernels loop' construct that has a 'num_workers' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_workers(1) worker(num:I) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{'num_workers' clause not allowed on a 'kernels loop' construct that has a 'worker' clause with an argument}} + // expected-note@+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@+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@-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@+4{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}} + // expected-note@-3{{previous clause is here}} + // expected-error@+2{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}} + // expected-note@-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@+4{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}} + // expected-note@-3{{previous clause is here}} + // expected-error@+2{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}} + // expected-note@-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@+4{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}} + // expected-note@-3{{previous clause is here}} + // expected-error@+2{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}} + // expected-note@-5{{previous clause is here}} +#pragma acc loop gang, worker, vector + for(int i = 0; i < 5; ++i); + } +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
