diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a8fd11a56edc83..61dd0c53581aa2 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12762,9 +12762,9 @@ def err_acc_gang_dim_value : Error<"argument to 'gang' clause dimension must be %select{a constant " "expression|1, 2, or 3: evaluated to %1}0">; def err_acc_num_arg_conflict - : Error<"'num' argument to '%0' clause not allowed on a '%1' " - "construct%select{| associated with a '%3' construct}2 that has a " - "'%4' clause">; + : Error<"'%0' argument to '%1' clause not allowed on a '%2' " + "construct%select{| associated with a '%4' construct}3 that has a " + "'%5' clause">; def err_acc_num_arg_conflict_reverse : Error<"'%0' clause not allowed on a 'kernels loop' construct that " "has a '%1' clause with a%select{n| 'num'}2 argument">; diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 5fc45974fa0ddc..18d58a43d265cf 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -804,6 +804,25 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause( if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) return nullptr; + // OpenACC 3.3 Section 2.9.4: + // An argument is allowed only when the 'vector_length' does not appear on the + // 'kernels' construct. + if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) { + auto VectorClauses = llvm::make_filter_range( + ExistingClauses, llvm::IsaPred); + + for (auto *VC : VectorClauses) { + if (cast(VC)->hasIntExpr()) { + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_num_arg_conflict_reverse) + << OpenACCClauseKind::VectorLength << OpenACCClauseKind::Vector + << /*num argument*/ 0; + SemaRef.Diag(VC->getBeginLoc(), diag::note_acc_previous_clause_here); + return nullptr; + } + } + } + assert(Clause.getIntExprs().size() == 1 && "Invalid number of expressions for NumWorkers"); return OpenACCVectorLengthClause::Create( @@ -1097,6 +1116,14 @@ ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, OpenACCGangKind GK, << HasAssocKind(DK, AssocKind) << AssocKind; return ExprError(); } +ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, StringRef TagKind, + OpenACCClauseKind CK, OpenACCDirectiveKind DK, + OpenACCDirectiveKind AssocKind) { + S.Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid) + << TagKind << CK << IsOrphanLoop(DK, AssocKind) << DK + << HasAssocKind(DK, AssocKind) << AssocKind; + return ExprError(); +} ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind, @@ -1172,8 +1199,9 @@ ExprResult CheckGangKernelsExpr(SemaOpenACC &S, if (Itr != Collection.end()) { S.Diag(E->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Gang << DK << HasAssocKind(DK, AssocKind) - << AssocKind << OpenACCClauseKind::NumGangs; + << "num" << OpenACCClauseKind::Gang << DK + << HasAssocKind(DK, AssocKind) << AssocKind + << OpenACCClauseKind::NumGangs; S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); return ExprError(); @@ -1206,63 +1234,94 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause( SemaOpenACC::OpenACCParsedClause &Clause) { 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: - // No restriction on when 'parallel' can contain an argument. - break; - case OpenACCDirectiveKind::Serial: - // GCC disallows this, and there is no real good reason for us to permit - // it, so disallow until we come up with a use case that makes sense. - DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num, - OpenACCClauseKind::Vector, Clause.getDirectiveKind(), - SemaRef.getActiveComputeConstructInfo().Kind); - IntExpr = nullptr; - break; - case OpenACCDirectiveKind::Kernels: { - const auto *Itr = - llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses, - llvm::IsaPred); - if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { - SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Vector << Clause.getDirectiveKind() - << HasAssocKind(Clause.getDirectiveKind(), - SemaRef.getActiveComputeConstructInfo().Kind) - << SemaRef.getActiveComputeConstructInfo().Kind - << OpenACCClauseKind::VectorLength; - SemaRef.Diag((*Itr)->getBeginLoc(), - diag::note_acc_previous_clause_here); + if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { + switch (SemaRef.getActiveComputeConstructInfo().Kind) { + case OpenACCDirectiveKind::Invalid: + case OpenACCDirectiveKind::Parallel: + // No restriction on when 'parallel' can contain an argument. + break; + case OpenACCDirectiveKind::Serial: + // GCC disallows this, and there is no real good reason for us to permit + // it, so disallow until we come up with a use case that makes sense. + DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector, + Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind); + IntExpr = nullptr; + break; + case OpenACCDirectiveKind::Kernels: { + const auto *Itr = + llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses, + llvm::IsaPred); + if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { + SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) + << "length" << OpenACCClauseKind::Vector + << Clause.getDirectiveKind() + << HasAssocKind(Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind) + << SemaRef.getActiveComputeConstructInfo().Kind + << OpenACCClauseKind::VectorLength; + 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::SerialLoop) { + DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector, + Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind); IntExpr = nullptr; + } else if (Clause.getDirectiveKind() == + OpenACCDirectiveKind::KernelsLoop) { + const auto *Itr = llvm::find_if( + ExistingClauses, llvm::IsaPred); + if (Itr != ExistingClauses.end()) { + SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) + << "length" << OpenACCClauseKind::Vector + << Clause.getDirectiveKind() + << HasAssocKind(Clause.getDirectiveKind(), + SemaRef.getActiveComputeConstructInfo().Kind) + << SemaRef.getActiveComputeConstructInfo().Kind + << OpenACCClauseKind::VectorLength; + 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.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::Vector << OpenACCClauseKind::Vector - << /*skip kernels construct info*/ 0; - SemaRef.Diag(SemaRef.LoopVectorClauseLoc, - diag::note_acc_previous_clause_here); - return nullptr; + if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { + // 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::Vector << OpenACCClauseKind::Vector + << /*skip kernels construct info*/ 0; + SemaRef.Diag(SemaRef.LoopVectorClauseLoc, + diag::note_acc_previous_clause_here); + return nullptr; + } } return OpenACCVectorClause::Create(Ctx, Clause.getBeginLoc(), @@ -1305,7 +1364,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause( llvm::IsaPred); if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Worker << Clause.getDirectiveKind() + << "num" << OpenACCClauseKind::Worker << Clause.getDirectiveKind() << HasAssocKind(Clause.getDirectiveKind(), SemaRef.getActiveComputeConstructInfo().Kind) << SemaRef.getActiveComputeConstructInfo().Kind @@ -1334,7 +1393,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause( llvm::IsaPred); if (Itr != ExistingClauses.end()) { SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict) - << OpenACCClauseKind::Worker << Clause.getDirectiveKind() + << "num" << OpenACCClauseKind::Worker << Clause.getDirectiveKind() << HasAssocKind(Clause.getDirectiveKind(), SemaRef.getActiveComputeConstructInfo().Kind) << SemaRef.getActiveComputeConstructInfo().Kind diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp index 45aaeb41a5ce0e..1a11f036b07aed 100644 --- a/clang/test/AST/ast-print-openacc-combined-construct.cpp +++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp @@ -344,4 +344,46 @@ void foo() { #pragma acc kernels loop worker(num:5) for(int i = 0;i<5;++i); + // CHECK: #pragma acc parallel loop vector +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop vector + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop vector(length: 5) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop vector(5) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop vector(length: 5) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop vector(length:5) + for(int i = 0;i<5;++i); + +// CHECK-NEXT: #pragma acc kernels loop vector +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc kernels loop vector + for(int i = 0;i<5;++i); + +// CHECK-NEXT: #pragma acc kernels loop vector(length: 5) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc kernels loop vector(5) + for(int i = 0;i<5;++i); + +// CHECK-NEXT: #pragma acc kernels loop vector(length: 5) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc kernels loop vector(length:5) + for(int i = 0;i<5;++i); + +// CHECK-NEXT: #pragma acc serial loop vector +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc serial loop vector + 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 a675db2a70a12f..45eede2e30f1f9 100644 --- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c @@ -48,7 +48,6 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop auto worker for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc parallel loop auto vector for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -167,7 +166,6 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop worker auto for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc parallel loop vector auto for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -287,7 +285,6 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop independent worker for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc parallel loop independent vector for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -406,7 +403,6 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop worker independent for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc parallel loop vector independent for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -650,9 +646,8 @@ void uses() { // 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}} - // TODOexpected-note@+1{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}} + // expected-note@+1{{previous clause is here}} #pragma acc parallel loop vector seq for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}} diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c index 3bf88d09c65e64..11bb342a2f638c 100644 --- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c @@ -39,7 +39,6 @@ void uses() { // 'worker', 'vector', 'seq', 'independent', 'auto', and 'tile' after // 'device_type'. - //expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} #pragma acc parallel loop device_type(*) vector for(int i = 0; i < 5; ++i); diff --git a/clang/test/SemaOpenACC/combined-construct-vector-ast.cpp b/clang/test/SemaOpenACC/combined-construct-vector-ast.cpp new file mode 100644 index 00000000000000..7b103021ccd555 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-vector-ast.cpp @@ -0,0 +1,176 @@ +// 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 vector + for(int i = 0; i < 5; ++i); + // CHECK: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop vector(Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop vector + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop vector + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop vector(Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +template +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 vector + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop vector(Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop vector(length:One) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop vector + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop vector + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop vector(Val) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop vector(length:One) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector 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: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral{{.*}}1 + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: ImplicitCastExpr + // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop + // CHECK-NEXT: vector clause + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral{{.*}}1 + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +void inst() { + TemplateUses(5); +} + +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/combined-construct-vector-clause.cpp b/clang/test/SemaOpenACC/combined-construct-vector-clause.cpp new file mode 100644 index 00000000000000..f96a36c1e395be --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-vector-clause.cpp @@ -0,0 +1,171 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +template +void TemplUses(T t) { + +#pragma acc parallel loop vector + for(int j = 0; j < 5; ++j); + +#pragma acc parallel loop vector(I) + for(int j = 0; j < 5; ++j); + +#pragma acc parallel loop vector(length:I) + for(int j = 0; j < 5; ++j); + +#pragma acc serial loop vector + for(int j = 0; j < 5; ++j); + + // expected-error@+1{{'length' argument on 'vector' clause is not permitted on a 'serial loop' construct}} +#pragma acc serial loop vector(I) + for(int j = 0; j < 5; ++j); + + // expected-error@+1{{'length' argument on 'vector' clause is not permitted on a 'serial loop' construct}} +#pragma acc serial loop vector(length:I) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector(I) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector(length:I) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector vector_length(t) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector_length(t) vector + for(int j = 0; j < 5; ++j); + + // expected-error@+2{{'vector_length' clause not allowed on a 'kernels loop' construct that has a 'vector' clause with an argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop vector(I) vector_length(t) + for(int j = 0; j < 5; ++j); + + // expected-error@+2{{'length' argument to 'vector' clause not allowed on a 'kernels loop' construct that has a 'vector_length' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop vector_length(t) vector(I) + for(int j = 0; j < 5; ++j); + +#pragma acc parallel loop vector + for(int j = 0; j < 5; ++j) { + // expected-error@+4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}} + // expected-error@+3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}} + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}} + // expected-note@-5 3{{previous clause is here}} +#pragma acc loop vector worker, gang + for(int j = 0; j < 5; ++j); + } +#pragma acc parallel loop vector + for(int j = 0; j < 5; ++j) { +#pragma acc serial loop vector worker, gang + for(int j = 0; j < 5; ++j); + } + +#pragma acc loop vector + for(int j = 0; j < 5; ++j) { +#pragma acc serial loop vector worker, gang + for(int j = 0; j < 5; ++j); + } + +#pragma acc kernels vector_length(t) + for(int j = 0; j < 5; ++j) { + // expected-error@+1{{'length' argument on 'vector' clause is not permitted on a 'serial loop' construct}} +#pragma acc serial loop vector(I) + for(int j = 0; j < 5; ++j); + } + +#pragma acc kernels vector_length(t) + for(int j = 0; j < 5; ++j) { +#pragma acc parallel loop vector(I) + for(int j = 0; j < 5; ++j); + } +} + +void uses() { + TemplUses(5); + + unsigned I; + int t; + +#pragma acc parallel loop vector + for(int j = 0; j < 5; ++j); + +#pragma acc parallel loop vector(I) + for(int j = 0; j < 5; ++j); + +#pragma acc parallel loop vector(length:I) + for(int j = 0; j < 5; ++j); + +#pragma acc serial loop vector + for(int j = 0; j < 5; ++j); + + // expected-error@+1{{'length' argument on 'vector' clause is not permitted on a 'serial loop' construct}} +#pragma acc serial loop vector(I) + for(int j = 0; j < 5; ++j); + + // expected-error@+1{{'length' argument on 'vector' clause is not permitted on a 'serial loop' construct}} +#pragma acc serial loop vector(length:I) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector(I) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector(length:I) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector vector_length(t) + for(int j = 0; j < 5; ++j); + +#pragma acc kernels loop vector_length(t) vector + for(int j = 0; j < 5; ++j); + + // expected-error@+2{{'vector_length' clause not allowed on a 'kernels loop' construct that has a 'vector' clause with an argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop vector(I) vector_length(t) + for(int j = 0; j < 5; ++j); + + // expected-error@+2{{'length' argument to 'vector' clause not allowed on a 'kernels loop' construct that has a 'vector_length' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop vector_length(t) vector(I) + for(int j = 0; j < 5; ++j); + +#pragma acc parallel loop vector + for(int j = 0; j < 5; ++j) { + // expected-error@+4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}} + // expected-error@+3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}} + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}} + // expected-note@-5 3{{previous clause is here}} +#pragma acc loop vector worker, gang + for(int j = 0; j < 5; ++j); + } +#pragma acc parallel loop vector + for(int j = 0; j < 5; ++j) { +#pragma acc serial loop vector worker, gang + for(int j = 0; j < 5; ++j); + } + +#pragma acc loop vector + for(int j = 0; j < 5; ++j) { +#pragma acc serial loop vector worker, gang + for(int j = 0; j < 5; ++j); + } + +#pragma acc kernels vector_length(t) + for(int j = 0; j < 5; ++j) { +#pragma acc parallel loop vector(I) + for(int j = 0; j < 5; ++j); + } + +#pragma acc kernels vector_length(t) + for(int j = 0; j < 5; ++j) { + // expected-error@+1{{'length' argument on 'vector' clause is not permitted on a 'serial loop' construct}} +#pragma acc serial loop vector(I) + for(int j = 0; j < 5; ++j); + } +} + diff --git a/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp b/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp index 3260e368fbb424..1fed82c5102083 100644 --- a/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp +++ b/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp @@ -18,12 +18,12 @@ void TemplUses(Int I, NotInt NI, ConvertsToInt CTI) { #pragma acc loop vector(length: NI) for(int j = 0; j < 5; ++j); - // expected-error@+2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} + // expected-error@+2{{'length' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} #pragma acc serial #pragma acc loop vector(length: I) for(int j = 0; j < 5; ++j); - // expected-error@+3{{'num' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}} + // expected-error@+3{{'length' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}} // expected-note@+1{{previous clause is here}} #pragma acc kernels vector_length(I) #pragma acc loop vector(length: CTI) @@ -87,12 +87,12 @@ void uses() { #pragma acc loop vector(length: NI) for(int j = 0; j < 5; ++j); - // expected-error@+2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} + // expected-error@+2{{'length' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} #pragma acc serial #pragma acc loop vector(length: i) for(int j = 0; j < 5; ++j); - // expected-error@+3{{'num' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}} + // expected-error@+3{{'length' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}} // expected-note@+1{{previous clause is here}} #pragma acc kernels vector_length(i) #pragma acc loop vector(length: i)