[clang] 0aa7892 - [OpenACC] Implement 'tile' for combined constructs
via cfe-commits
cfe-commits at lists.llvm.org
Mon Dec 2 13:31:42 PST 2024
Author: erichkeane
Date: 2024-12-02T13:31:36-08:00
New Revision: 0aa7892772a94b0395cb5a7fce077ab2348ecaa5
URL: https://github.com/llvm/llvm-project/commit/0aa7892772a94b0395cb5a7fce077ab2348ecaa5
DIFF: https://github.com/llvm/llvm-project/commit/0aa7892772a94b0395cb5a7fce077ab2348ecaa5.diff
LOG: [OpenACC] Implement 'tile' for combined constructs
Like collapse, this clause has only mild changes that need to be made.
Most of the associated work for the RAII container was already done
previously, so this is mostly just updating the diagnostics to properly
emit the right construct.
Added:
clang/test/SemaOpenACC/combined-construct-tile-ast.cpp
clang/test/SemaOpenACC/combined-construct-tile-clause.cpp
Modified:
clang/include/clang/Sema/SemaOpenACC.h
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
Removed:
################################################################################
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index b968ff779c89b0..d720cf3c74d87e 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -107,6 +107,10 @@ class SemaOpenACC : public SemaBase {
/// which allows us to diagnose if the number of arguments is too large for
/// the current number of 'for' loops.
bool TileDepthSatisfied = true;
+
+ /// Records the kind of the directive that this clause is attached to, which
+ /// allows us to use it in diagnostics.
+ OpenACCDirectiveKind DirectiveKind = OpenACCDirectiveKind::Invalid;
} TileInfo;
/// A list of the active reduction clauses, which allows us to check that all
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 7585bfca059567..654f3cd97c1c5e 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -598,9 +598,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // TODO OpenACC: Remove this when we get combined construct impl for this.
- if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
- return isNotImplemented();
// Duplicates here are not really sensible. We could possible permit
// multiples if they all had the same value, but there isn't really a good
@@ -1718,6 +1715,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
SemaRef.TileInfo.ActiveTile = TileClause;
SemaRef.TileInfo.TileDepthSatisfied = false;
SemaRef.TileInfo.CurTileCount = TileClause->getSizeExprs().size();
+ SemaRef.TileInfo.DirectiveKind = DirKind;
}
SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
@@ -2608,7 +2606,7 @@ void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
Diag(WhileLoc, diag::err_acc_invalid_in_loop)
- << /*while loop*/ 1 << OpenACCDirectiveKind::Loop
+ << /*while loop*/ 1 << TileInfo.DirectiveKind
<< OpenACCClauseKind::Tile;
assert(TileInfo.ActiveTile && "tile count without object?");
Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
@@ -2643,8 +2641,7 @@ void SemaOpenACC::ActOnDoStmt(SourceLocation DoLoc) {
if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
Diag(DoLoc, diag::err_acc_invalid_in_loop)
- << /*do loop*/ 2 << OpenACCDirectiveKind::Loop
- << OpenACCClauseKind::Tile;
+ << /*do loop*/ 2 << TileInfo.DirectiveKind << OpenACCClauseKind::Tile;
assert(TileInfo.ActiveTile && "tile count without object?");
Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
<< OpenACCClauseKind::Tile;
@@ -2692,7 +2689,7 @@ void SemaOpenACC::ForStmtBeginHelper(SourceLocation ForLoc,
if (LoopInfo.CurLevelHasLoopAlready) {
Diag(ForLoc, diag::err_acc_clause_multiple_loops)
- << OpenACCDirectiveKind::Loop << OpenACCClauseKind::Tile;
+ << TileInfo.DirectiveKind << OpenACCClauseKind::Tile;
assert(TileInfo.ActiveTile && "No tile object?");
Diag(TileInfo.ActiveTile->getBeginLoc(),
diag::note_acc_active_clause_here)
@@ -3203,7 +3200,7 @@ void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) {
if (OtherStmtLoc.isValid() && IsActiveTile) {
Diag(OtherStmtLoc, diag::err_acc_intervening_code)
- << OpenACCClauseKind::Tile << OpenACCDirectiveKind::Loop;
+ << OpenACCClauseKind::Tile << TileInfo.DirectiveKind;
Diag(TileInfo.ActiveTile->getBeginLoc(),
diag::note_acc_active_clause_here)
<< OpenACCClauseKind::Tile;
@@ -3232,7 +3229,7 @@ bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
}
if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
Diag(StartLoc, diag::err_acc_invalid_in_loop)
- << /*OpenACC Construct*/ 0 << OpenACCDirectiveKind::Loop
+ << /*OpenACC Construct*/ 0 << TileInfo.DirectiveKind
<< OpenACCClauseKind::Tile << K;
assert(TileInfo.ActiveTile && "Tile count without object?");
Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index 44ec0bb282f00e..d16e446706807a 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -1,5 +1,6 @@
// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
+constexpr int get_value() { return 1; }
void foo() {
int *iPtr;
// CHECK: #pragma acc parallel loop
@@ -210,4 +211,17 @@ void foo() {
#pragma acc parallel loop collapse(force:2)
for(int i = 0;i<5;++i)
for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop tile(1, 3, *, get_value())
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop tile(1, 3, *, get_value())
+ for(int i = 0;i<5;++i)
+ for(int i = 0;i<5;++i)
+ for(int i = 0;i<5;++i)
+ 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 9a12c1870e7379..fc5250ce548e44 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -158,7 +158,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto async
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
#pragma acc parallel loop auto tile(1+2, 1)
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
@@ -286,7 +285,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop async auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
#pragma acc parallel loop tile(1+2, 1) auto
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
@@ -415,7 +413,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent async
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
#pragma acc parallel loop independent tile(1+2, 1)
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
@@ -543,7 +540,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop async independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
#pragma acc parallel loop tile(1+2, 1) independent
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
@@ -678,7 +674,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop seq async
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
#pragma acc parallel loop seq tile(1+2, 1)
for(;;)
for(unsigned i = 0; i < 5; ++i);
@@ -812,7 +807,6 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop async seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
#pragma acc parallel loop tile(1+2, 1) seq
for(;;)
for(unsigned i = 0; i < 5; ++i);
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index 8933112a0063de..a5ab39cb12c383 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -209,7 +209,6 @@ void uses() {
#pragma acc parallel loop device_type(*) async
for(int i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc serial loop device_type(*) tile(*, 1)
for(int j = 0; j < 5; ++j)
for(int i = 0; i < 5; ++i);
diff --git a/clang/test/SemaOpenACC/combined-construct-tile-ast.cpp b/clang/test/SemaOpenACC/combined-construct-tile-ast.cpp
new file mode 100644
index 00000000000000..55a334c276ec8e
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-tile-ast.cpp
@@ -0,0 +1,327 @@
+// 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
+
+struct S {
+ constexpr S(){};
+ constexpr operator auto() {return 1;}
+};
+
+void NormalUses() {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop tile(S{}, 1, 2, *)
+ for(int i = 0;i < 5;++i)
+ for(int j = 0;j < 5;++j)
+ for(int k = 0;k < 5;++k)
+ for(int l = 0;l < 5;++l);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: tile clause
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 2
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} i 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} j 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} k 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'k' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'k' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} l 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'l' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'l' 'int'
+ // CHECK-NEXT: NullStmt
+}
+
+template<typename T, unsigned Value>
+void TemplUses() {
+ // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+ // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 Value
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()'
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc serial loop tile(S{}, T{}, *, T{} + S{}, Value, Value + 3)
+ for(int i = 0;i < 5;++i)
+ for(int j = 0;j < 5;++j)
+ for(int k = 0;k < 5;++k)
+ for(int l = 0;l < 5;++l)
+ for(int m = 0;m < 5;++m)
+ for(int n = 0;n < 5;++n);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop
+ // CHECK-NEXT: tile clause
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ //
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list
+ // CHECK-NEXT: InitListExpr{{.*}}'void'
+ //
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int'
+ //
+ // CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '+'
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list
+ // CHECK-NEXT: InitListExpr{{.*}}'void'
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ //
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'Value' 'unsigned int'
+ //
+ // CHECK-NEXT: BinaryOperator{{.*}}'unsigned int' '+'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'Value' 'unsigned int'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'unsigned int' <IntegralCast>
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} i 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} j 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} k 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'k' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'k' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} l 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'l' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'l' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} m 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'m' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'m' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} n 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'n' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'n' 'int'
+ // CHECK-NEXT: NullStmt
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void ()' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'S'
+ // CHECK-NEXT: RecordType{{.*}} 'S'
+ // CHECK-NEXT: CXXRecord{{.*}} 'S'
+ // CHECK-NEXT: TemplateArgument integral '2U'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop
+ // CHECK-NEXT: tile clause
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ //
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int'
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 2
+ // CHECK-NEXT: BinaryOperator{{.*}}'int' '+'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 2
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 Value
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 2
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 5
+ // CHECK-NEXT: BinaryOperator{{.*}}'unsigned int' '+'
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 Value
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 2
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'unsigned int' <IntegralCast>
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3
+
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} i 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} j 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} k 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'k' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'k' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} l 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'l' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'l' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} m 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'m' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'m' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} n 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: BinaryOperator{{.*}}'<'
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'n' 'int'
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}}++
+ // CHECK-NEXT: DeclRefExpr{{.*}}'n' 'int'
+ // CHECK-NEXT: NullStmt
+
+}
+
+void Inst() {
+ TemplUses<S, 2>();
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp b/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp
new file mode 100644
index 00000000000000..00c551e163666a
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp
@@ -0,0 +1,396 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+constexpr int three() { return 3; }
+constexpr int one() { return 1; }
+constexpr int neg() { return -1; }
+constexpr int zero() { return 0; }
+
+struct NotConstexpr {
+ constexpr NotConstexpr(){};
+
+ operator int(){ return 1; }
+};
+struct ConvertsNegative {
+ constexpr ConvertsNegative(){};
+
+ constexpr operator int(){ return -1; }
+};
+struct ConvertsOne{
+ constexpr ConvertsOne(){};
+
+ constexpr operator int(){ return 1; }
+};
+
+struct ConvertsThree{
+ constexpr ConvertsThree(){};
+
+ constexpr operator int(){ return 3; }
+};
+
+template<typename T, int Val>
+void negative_zero_constexpr_templ() {
+ // expected-error at +1 2{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc serial loop tile(*, T{})
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc parallel loop tile(Val, *)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc kernels loop tile(zero(), *)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+}
+
+void negative_zero_constexpr() {
+ negative_zero_constexpr_templ<int, 1>(); // expected-note{{in instantiation of function template specialization}}
+ negative_zero_constexpr_templ<int, -1>(); // expected-note{{in instantiation of function template specialization}}
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc serial loop tile(0, *)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc parallel loop tile(1, 0)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc kernels loop tile(1, -1)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc parallel loop tile(-1, 0)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc serial loop tile(zero(), 0)
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc kernels loop tile(1, neg())
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be an asterisk or a constant expression}}
+#pragma acc parallel loop tile(NotConstexpr{})
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc serial loop tile(1, ConvertsNegative{})
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+#pragma acc kernels loop tile(*, ConvertsOne{})
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+}
+
+template<unsigned One>
+void only_for_loops_templ() {
+ // expected-note at +1{{'parallel loop' construct is here}}
+#pragma acc parallel loop tile(One)
+ // expected-error at +1{{OpenACC 'parallel loop' construct can only be applied to a 'for' loop}}
+ while(true);
+
+ // expected-note at +1{{'serial loop' construct is here}}
+#pragma acc serial loop tile(One)
+ // expected-error at +1{{OpenACC 'serial loop' construct can only be applied to a 'for' loop}}
+ do {} while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc kernels loop tile(One, 2) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ // expected-error at +1{{while loop cannot appear in intervening code of a 'kernels loop' with a 'tile' clause}}
+ while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc serial loop tile(One, 2) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ // expected-error at +1{{do loop cannot appear in intervening code of a 'serial loop' with a 'tile' clause}}
+ do{}while(true);
+}
+
+
+void only_for_loops() {
+ // expected-note at +1{{'parallel loop' construct is here}}
+#pragma acc parallel loop tile(1)
+ // expected-error at +1{{OpenACC 'parallel loop' construct can only be applied to a 'for' loop}}
+ while(true);
+
+ // expected-note at +1{{'serial loop' construct is here}}
+#pragma acc serial loop tile(1)
+ // expected-error at +1{{OpenACC 'serial loop' construct can only be applied to a 'for' loop}}
+ do {} while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc kernels loop tile(1, 2) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ // expected-error at +1{{while loop cannot appear in intervening code of a 'kernels loop' with a 'tile' clause}}
+ while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc parallel loop tile(1, 2) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ // expected-error at +1{{do loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+ do{}while(true);
+}
+
+void only_one_on_loop() {
+ // expected-error at +2{{OpenACC 'tile' clause cannot appear more than once on a 'serial loop' directive}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc serial loop tile(1) tile(1)
+ for(int i = 0; i < 5; ++i);
+}
+
+template<unsigned Val>
+void depth_too_high_templ() {
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc kernels loop tile (Val, *, Val) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc parallel loop tile (Val, *, Val) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j)
+ // expected-error at +1{{while loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+ while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc serial loop tile (Val, *, Val) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j)
+ // expected-error at +1{{do loop cannot appear in intervening code of a 'serial loop' with a 'tile' clause}}
+ do{}while(true);
+
+ int Arr[Val+5];
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc kernels loop tile (Val, *, Val) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(auto x : Arr)
+ // expected-error at +1{{while loop cannot appear in intervening code of a 'kernels loop' with a 'tile' clause}}
+ while(true)
+ for(int j = 0; j < 5; ++j);
+
+#pragma acc parallel loop tile (Val, *, Val)
+ for(int i = 0; i < 5; ++i)
+ for(auto x : Arr)
+ for(int j = 0; j < 5; ++j)
+ while(true);
+}
+
+void depth_too_high() {
+ depth_too_high_templ<3>();
+
+int Arr[5];
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc serial loop tile (1, *, 3) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc parallel loop tile (1, *, 3) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j)
+ // expected-error at +1{{while loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+ while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc parallel loop tile (1, *, 3) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j)
+ // expected-error at +1{{do loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+ do{}while(true);
+
+ // expected-error at +1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc parallel loop tile (1, *, 3) // expected-note 2{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i)
+ for(int j = 0; j < 5; ++j)
+ // expected-error at +1{{while loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+ while(true)
+ for(int j = 0; j < 5; ++j);
+
+#pragma acc parallel loop tile (1, *, 3)
+ for(int i = 0; i < 5; ++i)
+ for(auto x : Arr)
+ for(int j = 0; j < 5; ++j)
+ while(true);
+}
+
+template<unsigned Val>
+void not_single_loop_templ() {
+
+ int Arr[Val];
+
+#pragma acc parallel loop tile (Val, *, 3) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ for (auto x : Arr)
+ for(int k = 0; k < 5; ++k);
+ // expected-error at +1{{more than one for-loop in a loop associated with OpenACC 'parallel loop' construct with a 'tile' clause}}
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k);
+ }
+}
+
+void not_single_loop() {
+ not_single_loop_templ<3>(); // no diagnostic, was diagnosed in phase 1.
+
+ int Arr[5];
+
+#pragma acc parallel loop tile (1, *, 3)// expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ for (auto x : Arr)
+ for(int k = 0; k < 5; ++k);
+ // expected-error at +1{{more than one for-loop in a loop associated with OpenACC 'parallel loop' construct with a 'tile' clause}}
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k);
+ }
+}
+
+template<unsigned Val>
+void no_other_directives_templ() {
+
+ int Arr[Val];
+
+#pragma acc parallel loop tile (Val, *, 3) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ for (auto x : Arr) {
+ // expected-error at +1{{OpenACC 'serial' construct cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+#pragma acc serial
+ ;
+ for(int j = 0; j < 5; ++j);
+ }
+ }
+
+ // OK, in innermost
+#pragma acc parallel loop tile (Val, *, 3)
+ for(int i = 0; i < 5; ++i) {
+ for(int j = 0; j < 5; ++j) {
+ for (auto x : Arr) {
+#pragma acc serial
+ ;
+ }
+ }
+ }
+}
+
+void no_other_directives() {
+ no_other_directives_templ<3>();
+ int Arr[5];
+
+#pragma acc parallel loop tile (1, *, 3) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ for (auto x : Arr) {
+ // expected-error at +1{{OpenACC 'serial' construct cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}}
+#pragma acc serial
+ ;
+ for(int j = 0; j < 5; ++j);
+ }
+ }
+
+ // OK, in innermost
+#pragma acc parallel loop tile (3, *, 3)
+ for(int i = 0; i < 5; ++i) {
+ for(int j = 0; j < 5; ++j) {
+ for (auto x : Arr) {
+#pragma acc serial
+ ;
+ }
+ }
+ }
+}
+
+void call();
+template<unsigned Val>
+void intervening_templ() {
+#pragma acc parallel loop tile(1, Val, *) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ //expected-error at +1{{inner loops must be tightly nested inside a 'tile' clause on a 'parallel loop' construct}}
+ call();
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k);
+ }
+
+#pragma acc parallel loop tile(1, Val, *) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ //expected-error at +1{{inner loops must be tightly nested inside a 'tile' clause on a 'parallel loop' construct}}
+ unsigned I;
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k);
+ }
+
+#pragma acc parallel loop tile(1, Val, *)
+ // expected-error at +2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+ // expected-note at -2{{'parallel loop' construct is here}}
+ for(int i = 0;;++i) {
+ // expected-error at +2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+ // expected-note at -5{{'parallel loop' construct is here}}
+ for(int j = 0;;++j)
+ // expected-error at +2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+ // expected-note at -8{{'parallel loop' construct is here}}
+ for(int k = 0;;++k)
+ call();
+ }
+}
+
+void intervening() {
+ intervening_templ<3>();
+
+#pragma acc parallel loop tile(1, 2, *) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ //expected-error at +1{{inner loops must be tightly nested inside a 'tile' clause on a 'parallel loop' construct}}
+ call();
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k);
+ }
+
+#pragma acc parallel loop tile(1, 2, *) // expected-note{{active 'tile' clause defined here}}
+ for(int i = 0; i < 5; ++i) {
+ //expected-error at +1{{inner loops must be tightly nested inside a 'tile' clause on a 'parallel loop' construct}}
+ unsigned I;
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k);
+ }
+
+#pragma acc parallel loop tile(1, 2, *)
+ for(int i = 0; i < 5; ++i) {
+ for(int j = 0; j < 5; ++j)
+ for(int k = 0; k < 5; ++k)
+ call();
+ }
+
+#pragma acc parallel loop tile(1, 2, *)
+ // expected-error at +2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+ // expected-note at -2{{'parallel loop' construct is here}}
+ for(int i = 0;;++i) {
+ // expected-error at +2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+ // expected-note at -5{{'parallel loop' construct is here}}
+ for(int j = 0;;++j)
+ // expected-error at +2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+ // expected-note at -8{{'parallel loop' construct is here}}
+ for(int k = 0;;++k)
+ for(;;)
+ call();
+ }
+}
+
+void collapse_tile_depth() {
+ // expected-error at +4{{'collapse' clause specifies a loop count greater than the number of available loops}}
+ // expected-note at +3{{active 'collapse' clause defined here}}
+ // expected-error at +2{{'tile' clause specifies a loop count greater than the number of available loops}}
+ // expected-note at +1{{active 'tile' clause defined here}}
+#pragma acc parallel loop tile(1, 2, 3) collapse (3)
+ for(int i = 0; i < 5;++i) {
+ for(int j = 0; j < 5; ++j);
+ }
+}
More information about the cfe-commits
mailing list