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@+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' <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@+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<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@+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<unsigned One> +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<unsigned Val> +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<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@+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<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@+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<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@+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); + } +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
