[clang] 009b5e8 - [OpenACC] 'vector' clause implementation for combined constructs
via cfe-commits
cfe-commits at lists.llvm.org
Fri Dec 6 09:07:03 PST 2024
Author: erichkeane
Date: 2024-12-06T09:06:57-08:00
New Revision: 009b5e8e5915910d172f6660ceb69784c18e7ac7
URL: https://github.com/llvm/llvm-project/commit/009b5e8e5915910d172f6660ceb69784c18e7ac7
DIFF: https://github.com/llvm/llvm-project/commit/009b5e8e5915910d172f6660ceb69784c18e7ac7.diff
LOG: [OpenACC] 'vector' clause implementation for combined constructs
Similar to 'worker', the 'vector' clause has some rules that needed to
be applied on its argument legality that for combined constructs need to
look at the current construct, not the 'effective' parent construct.
Additionally, it has some interaction with `vector_length` that needed
to be encoded as well. This patch implements it.
Added:
clang/test/SemaOpenACC/combined-construct-vector-ast.cpp
clang/test/SemaOpenACC/combined-construct-vector-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/loop-construct-vector-clause.cpp
Removed:
################################################################################
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<OpenACCVectorClause>);
+
+ for (auto *VC : VectorClauses) {
+ if (cast<OpenACCVectorClause>(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<OpenACCVectorLengthClause>);
- 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<OpenACCVectorLengthClause>);
+ 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<OpenACCVectorLengthClause>);
+ 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<OpenACCNumWorkersClause>);
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<OpenACCNumWorkersClause>);
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 at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop auto vector
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +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 at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop vector auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +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 at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop independent vector
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +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 at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc parallel loop vector independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'nohost' not yet implemented}}
@@ -650,9 +646,8 @@ void uses() {
// 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}}
- // TODOexpected-note at +1{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc parallel loop vector seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +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 at +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<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 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<int, 1>(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<typename T, unsigned I>
+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 at +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 at +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 at +2{{'vector_length' clause not allowed on a 'kernels loop' construct that has a 'vector' clause with an argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop vector(I) vector_length(t)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +2{{'length' argument to 'vector' clause not allowed on a 'kernels loop' construct that has a 'vector_length' clause}}
+ // expected-note at +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 at +4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}}
+ // expected-note at -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 at +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<int, 5>(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 at +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 at +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 at +2{{'vector_length' clause not allowed on a 'kernels loop' construct that has a 'vector' clause with an argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop vector(I) vector_length(t)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +2{{'length' argument to 'vector' clause not allowed on a 'kernels loop' construct that has a 'vector_length' clause}}
+ // expected-note at +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 at +4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}}
+ // expected-note at -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 at +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 at +2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+ // expected-error at +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 at +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 at +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 at +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 at +2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+ // expected-error at +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 at +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 at +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 at +1{{previous clause is here}}
#pragma acc kernels vector_length(i)
#pragma acc loop vector(length: i)
More information about the cfe-commits
mailing list