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@+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@+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@+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@+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@+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@+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@+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' + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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 +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' + // 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{{.*}}'' '+' + // 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' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3 + // CHECK-NEXT: ForStmt + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} i 'int' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0 + // CHECK-NEXT: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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' + // 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' + // 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' + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' + // 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' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3 + + // CHECK-NEXT: ForStmt + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} i 'int' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0 + // CHECK-NEXT: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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: <<>> + // 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(); +} +#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 +void negative_zero_constexpr_templ() { + // expected-error@+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@+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@+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(); // expected-note{{in instantiation of function template specialization}} + negative_zero_constexpr_templ(); // expected-note{{in instantiation of function template specialization}} + + // expected-error@+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@+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@+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@+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@+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@+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@+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@+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 +void only_for_loops_templ() { + // expected-note@+1{{'parallel loop' construct is here}} +#pragma acc parallel loop tile(One) + // expected-error@+1{{OpenACC 'parallel loop' construct can only be applied to a 'for' loop}} + while(true); + + // expected-note@+1{{'serial loop' construct is here}} +#pragma acc serial loop tile(One) + // expected-error@+1{{OpenACC 'serial loop' construct can only be applied to a 'for' loop}} + do {} while(true); + + // expected-error@+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@+1{{while loop cannot appear in intervening code of a 'kernels loop' with a 'tile' clause}} + while(true); + + // expected-error@+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@+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@+1{{'parallel loop' construct is here}} +#pragma acc parallel loop tile(1) + // expected-error@+1{{OpenACC 'parallel loop' construct can only be applied to a 'for' loop}} + while(true); + + // expected-note@+1{{'serial loop' construct is here}} +#pragma acc serial loop tile(1) + // expected-error@+1{{OpenACC 'serial loop' construct can only be applied to a 'for' loop}} + do {} while(true); + + // expected-error@+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@+1{{while loop cannot appear in intervening code of a 'kernels loop' with a 'tile' clause}} + while(true); + + // expected-error@+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@+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@+2{{OpenACC 'tile' clause cannot appear more than once on a 'serial loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc serial loop tile(1) tile(1) + for(int i = 0; i < 5; ++i); +} + +template +void depth_too_high_templ() { + // expected-error@+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@+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@+1{{while loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}} + while(true); + + // expected-error@+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@+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@+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@+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@+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@+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@+1{{while loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}} + while(true); + + // expected-error@+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@+1{{do loop cannot appear in intervening code of a 'parallel loop' with a 'tile' clause}} + do{}while(true); + + // expected-error@+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@+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 +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@+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@+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 +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@+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@+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 +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@+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@+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@+2{{OpenACC 'parallel loop' construct must have a terminating condition}} + // expected-note@-2{{'parallel loop' construct is here}} + for(int i = 0;;++i) { + // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}} + // expected-note@-5{{'parallel loop' construct is here}} + for(int j = 0;;++j) + // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}} + // expected-note@-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@+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@+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@+2{{OpenACC 'parallel loop' construct must have a terminating condition}} + // expected-note@-2{{'parallel loop' construct is here}} + for(int i = 0;;++i) { + // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}} + // expected-note@-5{{'parallel loop' construct is here}} + for(int j = 0;;++j) + // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}} + // expected-note@-8{{'parallel loop' construct is here}} + for(int k = 0;;++k) + for(;;) + call(); + } +} + +void collapse_tile_depth() { + // expected-error@+4{{'collapse' clause specifies a loop count greater than the number of available loops}} + // expected-note@+3{{active 'collapse' clause defined here}} + // expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}} + // expected-note@+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); + } +}