Author: erichkeane Date: 2024-12-03T09:31:40-08:00 New Revision: eb257fe37ba1ea41bd162e2fbd0ee4cd33fcdcea
URL: https://github.com/llvm/llvm-project/commit/eb257fe37ba1ea41bd162e2fbd0ee4cd33fcdcea DIFF: https://github.com/llvm/llvm-project/commit/eb257fe37ba1ea41bd162e2fbd0ee4cd33fcdcea.diff LOG: [OpenACC] Enable 3 more clauses for combined constructs. 'num_gangs', 'num_workers', and 'vector_length' are similar to eachother, and are all the same implementation as for compute constructs, so this patch just enables them and adds the necessary testing to ensure they work correctly. These will get more complicated when they get combined with 'gang', 'worker', 'vector' and 'reduction', but those restrictions will be implemented when those clauses are enabled. Added: clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp clang/test/SemaOpenACC/combined-construct-num_workers-clause.c clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp clang/test/SemaOpenACC/combined-construct-vector_length-clause.c Modified: 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/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 654f3cd97c1c5e..060df967322ac5 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -694,14 +694,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions only properly implemented on 'compute' constructs, and - // 'compute' constructs are the only construct that can do anything with - // this yet, so skip/treat as unimplemented in this case. - // TODO OpenACC: Remove this check when we have combined constructs for this - // clause. - if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind())) - return isNotImplemented(); - // There is no prose in the standard that says duplicates aren't allowed, // but this diagnostic is present in other compilers, as well as makes // sense. @@ -730,6 +722,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( // OpenACC 3.3 Section 2.5.4: // A reduction clause may not appear on a parallel construct with a // num_gangs clause that has more than one argument. + // TODO: OpenACC: Reduction on Combined Construct needs to do this too. if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel && Clause.getIntExprs().size() > 1) { auto *Parallel = @@ -751,13 +744,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions only properly implemented on 'compute' constructs, and - // 'compute' constructs are the only construct that can do anything with - // this yet, so skip/treat as unimplemented in this case. - // TODO: OpenACC: Remove when we get combined constructs. - if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind())) - return isNotImplemented(); - // There is no prose in the standard that says duplicates aren't allowed, // but this diagnostic is present in other compilers, as well as makes // sense. @@ -773,13 +759,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions only properly implemented on 'compute' constructs, and - // 'compute' constructs are the only construct that can do anything with - // this yet, so skip/treat as unimplemented in this case. - // TODO: OpenACC: Remove when we get combined constructs. - if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind())) - return isNotImplemented(); - // There is no prose in the standard that says duplicates aren't allowed, // but this diagnostic is present in other compilers, as well as makes // sense. diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp index d16e446706807a..435c770c7457d1 100644 --- a/clang/test/AST/ast-print-openacc-combined-construct.cpp +++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp @@ -224,4 +224,22 @@ void foo() { for(int i = 0;i<5;++i) for(int i = 0;i<5;++i); +// CHECK: #pragma acc parallel loop num_gangs(i, (int)array[2]) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop num_gangs(i, (int)array[2]) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop num_workers(i) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop num_workers(i) + for(int i = 0;i<5;++i); + +// CHECK: #pragma acc parallel loop vector_length((int)array[1]) +// CHECK-NEXT: for (int i = 0; i < 5; ++i) +// CHECK-NEXT: ; +#pragma acc parallel loop vector_length((int)array[1]) + 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 fc5250ce548e44..a6f57a63a91ddf 100644 --- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c @@ -134,16 +134,10 @@ void uses() { // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc parallel loop auto bind(Var) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}} #pragma acc parallel loop auto vector_length(1) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}} #pragma acc parallel loop auto num_gangs(1) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}} #pragma acc parallel loop auto num_workers(1) for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}} @@ -261,16 +255,10 @@ void uses() { // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc parallel loop bind(Var) auto for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}} #pragma acc parallel loop vector_length(1) auto for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}} #pragma acc parallel loop num_gangs(1) auto for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}} #pragma acc parallel loop num_workers(1) auto for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}} @@ -389,16 +377,10 @@ void uses() { // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc parallel loop independent bind(Var) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}} #pragma acc parallel loop independent vector_length(1) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}} #pragma acc parallel loop independent num_gangs(1) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}} #pragma acc parallel loop independent num_workers(1) for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}} @@ -516,16 +498,10 @@ void uses() { // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc parallel loop bind(Var) independent for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}} #pragma acc parallel loop vector_length(1) independent for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}} #pragma acc parallel loop num_gangs(1) independent for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}} #pragma acc parallel loop num_workers(1) independent for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}} @@ -650,16 +626,10 @@ void uses() { // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc parallel loop seq bind(Var) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}} #pragma acc parallel loop seq vector_length(1) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}} #pragma acc parallel loop seq num_gangs(1) for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}} #pragma acc parallel loop seq num_workers(1) for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}} @@ -783,16 +753,10 @@ void uses() { // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc parallel loop bind(Var) seq for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}} #pragma acc parallel loop vector_length(1) seq for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}} #pragma acc parallel loop num_gangs(1) seq for(unsigned i = 0; i < 5; ++i); - // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}} - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}} #pragma acc parallel loop num_workers(1) seq for(unsigned i = 0; i < 5; ++i); // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}} diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c index a5ab39cb12c383..9a60fb4c665e5a 100644 --- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c @@ -195,7 +195,6 @@ void uses() { // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial loop' directive}} #pragma acc serial loop device_type(*) num_gangs(1) for(int i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented, clause ignored}} #pragma acc parallel loop device_type(*) num_workers(1) for(int i = 0; i < 5; ++i); // expected-error@+2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a 'serial loop' construct}} diff --git a/clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp b/clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp new file mode 100644 index 00000000000000..6e75a009433645 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp @@ -0,0 +1,121 @@ +// 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 +int some_int(); +short some_short(); +long some_long(); +struct CorrectConvert { + operator int(); +} Convert; + + +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt + +#pragma acc parallel loop num_gangs(some_int(), some_long(), some_short()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: CallExpr{{.*}}'int' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()' + // CHECK-NEXT: CallExpr{{.*}}'long' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()' + // CHECK-NEXT: CallExpr{{.*}}'short' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop num_gangs(some_int()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: CallExpr{{.*}}'int' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +template<typename T, typename U> +void TemplUses(T t, U u) { + // CHECK-NEXT: FunctionTemplateDecl + // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T + // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U + // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)' + // CHECK-NEXT: ParmVarDecl{{.*}} t 'T' + // CHECK-NEXT: ParmVarDecl{{.*}} u 'U' + // CHECK-NEXT: CompoundStmt + +#pragma acc kernels loop num_gangs(u) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop num_gangs(u, U::value) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // Check the instantiated versions of the above. + // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'CorrectConvert' + // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert' + // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert' + // CHECK-NEXT: TemplateArgument type 'HasInt' + // CHECK-NEXT: RecordType{{.*}} 'HasInt' + // CHECK-NEXT: CXXRecord{{.*}} 'HasInt' + // CHECK-NEXT: ParmVarDecl{{.*}} t 'CorrectConvert' + // CHECK-NEXT: ParmVarDecl{{.*}} u 'HasInt' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +struct HasInt { + using IntTy = int; + using ShortTy = short; + static constexpr int value = 1; + + operator char(); +}; + +void Inst() { + TemplUses<CorrectConvert, HasInt>({}, {}); +} +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c new file mode 100644 index 00000000000000..bd035bd4a51a27 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c @@ -0,0 +1,45 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +short getS(); +float getF(); +void Test() { +#pragma acc kernels loop num_gangs(1) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial loop' directive}} +#pragma acc serial loop num_gangs(1) + for(int i = 5; i < 10;++i); + +#pragma acc parallel loop num_gangs(1) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type}} +#pragma acc parallel loop num_gangs(getF()) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{expected expression}} +#pragma acc kernels loop num_gangs() + for(int i = 5; i < 10;++i); + + // expected-error@+1{{expected expression}} +#pragma acc parallel loop num_gangs() + for(int i = 5; i < 10;++i); + + // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(1) num_gangs(2) + for(int i = 5; i < 10;++i); + + // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop num_gangs(1) num_gangs(2) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels loop' directive expects maximum of 1, 2 were provided}} +#pragma acc kernels loop num_gangs(1, getS()) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel loop' directive expects maximum of 3, 4 were provided}} +#pragma acc parallel loop num_gangs(getS(), 1, getS(), 1) + for(int i = 5; i < 10;++i); +} diff --git a/clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp b/clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp new file mode 100644 index 00000000000000..8aa361c7b037c0 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp @@ -0,0 +1,230 @@ +// 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 +int some_int(); +short some_short(); +long some_long(); +enum E{}; +E some_enum(); +struct CorrectConvert { + operator int(); +} Convert; + + +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt +#pragma acc parallel loop num_workers(some_int()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CallExpr{{.*}}'int' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop num_workers(some_short()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CallExpr{{.*}}'short' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop num_workers(some_long()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CallExpr{{.*}}'long' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop num_workers(some_enum()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CallExpr{{.*}}'E' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'E (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'E ()' lvalue Function{{.*}} 'some_enum' 'E ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop num_workers(Convert) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int + // CHECK-NEXT: DeclRefExpr{{.*}} 'struct CorrectConvert':'CorrectConvert' lvalue Var + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +template<typename T, typename U> +void TemplUses(T t, U u) { + // CHECK-NEXT: FunctionTemplateDecl + // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T + // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U + // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)' + // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T' + // CHECK-NEXT: ParmVarDecl{{.*}} referenced u 'U' + // CHECK-NEXT: CompoundStmt + +#pragma acc parallel loop num_workers(t) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue ParmVar{{.*}} 't' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop num_workers(u) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop num_workers(U::value) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop num_workers(T{}) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list + // CHECK-NEXT: InitListExpr{{.*}} 'void' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop num_workers(U{}) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'U' 'U' list + // CHECK-NEXT: InitListExpr{{.*}} 'void' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc kernels loop num_workers(typename U::IntTy{}) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename U::IntTy' 'typename U::IntTy' list + // CHECK-NEXT: InitListExpr{{.*}} 'void' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop num_workers(typename U::ShortTy{}) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename U::ShortTy' 'typename U::ShortTy' list + // CHECK-NEXT: InitListExpr{{.*}} 'void' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // Check the instantiated versions of the above. + // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'CorrectConvert' + // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert' + // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert' + // CHECK-NEXT: TemplateArgument type 'HasInt' + // CHECK-NEXT: RecordType{{.*}} 'HasInt' + // CHECK-NEXT: CXXRecord{{.*}} 'HasInt' + // CHECK-NEXT: ParmVarDecl{{.*}} used t 'CorrectConvert' + // CHECK-NEXT: ParmVarDecl{{.*}} used u 'HasInt' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int + // CHECK-NEXT: DeclRefExpr{{.*}} 'CorrectConvert' lvalue ParmVar + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'CorrectConvert' lvalue + // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'CorrectConvert' functional cast to struct CorrectConvert <NoOp> + // CHECK-NEXT: InitListExpr{{.*}}'CorrectConvert' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'HasInt' lvalue + // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'HasInt' functional cast to struct HasInt <NoOp> + // CHECK-NEXT: InitListExpr{{.*}}'HasInt' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'typename HasInt::IntTy':'int' functional cast to typename struct HasInt::IntTy <NoOp> + // CHECK-NEXT: InitListExpr{{.*}}'typename HasInt::IntTy':'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: num_workers clause + // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'typename HasInt::ShortTy':'short' functional cast to typename struct HasInt::ShortTy <NoOp> + // CHECK-NEXT: InitListExpr{{.*}}'typename HasInt::ShortTy':'short' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} +struct HasInt { + using IntTy = int; + using ShortTy = short; + static constexpr int value = 1; + + operator char(); +}; + +void Inst() { + TemplUses<CorrectConvert, HasInt>({}, {}); +} +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c b/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c new file mode 100644 index 00000000000000..a5891f071bb030 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c @@ -0,0 +1,37 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +short getS(); +float getF(); +void Test() { +#pragma acc kernels loop num_workers(1) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{OpenACC 'num_workers' clause is not valid on 'serial loop' directive}} +#pragma acc serial loop num_workers(1) + for(int i = 5; i < 10;++i); + +#pragma acc parallel loop num_workers(1) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{OpenACC clause 'num_workers' requires expression of integer type}} +#pragma acc parallel loop num_workers(getF()) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{expected expression}} +#pragma acc kernels loop num_workers() + for(int i = 5; i < 10;++i); + + // expected-error@+1{{expected expression}} +#pragma acc parallel loop num_workers() + for(int i = 5; i < 10;++i); + + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} +#pragma acc kernels loop num_workers(1, 2) + for(int i = 5; i < 10;++i); + + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} +#pragma acc parallel loop num_workers(1, 2) + for(int i = 5; i < 10;++i); +} diff --git a/clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp b/clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp new file mode 100644 index 00000000000000..6cfc9c6b8b2c25 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp @@ -0,0 +1,98 @@ +// 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 +short some_short(); + +struct CorrectConvert { + operator int(); +} Convert; + + +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt + +#pragma acc kernels loop vector_length(some_short()) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: CallExpr{{.*}}'short' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} +template<typename T, typename U> +void TemplUses(T t, U u) { + // CHECK-NEXT: FunctionTemplateDecl + // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T + // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U + // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)' + // CHECK-NEXT: ParmVarDecl{{.*}} t 'T' + // CHECK-NEXT: ParmVarDecl{{.*}} u 'U' + // CHECK-NEXT: CompoundStmt + +#pragma acc kernels loop vector_length(u) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc parallel loop vector_length(U::value) + for (unsigned i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // Check the instantiated versions of the above. + // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'CorrectConvert' + // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert' + // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert' + // CHECK-NEXT: TemplateArgument type 'HasInt' + // CHECK-NEXT: RecordType{{.*}} 'HasInt' + // CHECK-NEXT: CXXRecord{{.*}} 'HasInt' + // CHECK-NEXT: ParmVarDecl{{.*}} t 'CorrectConvert' + // CHECK-NEXT: ParmVarDecl{{.*}} u 'HasInt' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: vector_length clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt +} + +struct HasInt { + using IntTy = int; + using ShortTy = short; + static constexpr int value = 1; + + operator char(); +}; + +void Inst() { + TemplUses<CorrectConvert, HasInt>({}, {}); +} +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c b/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c new file mode 100644 index 00000000000000..8b6dedd9b83baa --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c @@ -0,0 +1,37 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +short getS(); +float getF(); +void Test() { +#pragma acc kernels loop vector_length(1) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial loop' directive}} +#pragma acc serial loop vector_length(1) + for(int i = 5; i < 10;++i); + +#pragma acc parallel loop vector_length(1) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type}} +#pragma acc parallel loop vector_length(getF()) + for(int i = 5; i < 10;++i); + + // expected-error@+1{{expected expression}} +#pragma acc kernels loop vector_length() + for(int i = 5; i < 10;++i); + + // expected-error@+1{{expected expression}} +#pragma acc parallel loop vector_length() + for(int i = 5; i < 10;++i); + + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} +#pragma acc kernels loop vector_length(1, 2) + for(int i = 5; i < 10;++i); + + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} +#pragma acc parallel loop vector_length(1, 2) + for(int i = 5; i < 10;++i); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
