[clang] eb257fe - [OpenACC] Enable 3 more clauses for combined constructs.
via cfe-commits
cfe-commits at lists.llvm.org
Tue Dec 3 09:48:30 PST 2024
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 at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc parallel loop auto bind(Var)
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC clause 'device_num' not yet implemented}}
@@ -261,16 +255,10 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc parallel loop bind(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC clause 'device_num' not yet implemented}}
@@ -389,16 +377,10 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc parallel loop independent bind(Var)
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC clause 'device_num' not yet implemented}}
@@ -516,16 +498,10 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc parallel loop bind(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC clause 'device_num' not yet implemented}}
@@ -650,16 +626,10 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc parallel loop seq bind(Var)
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC clause 'device_num' not yet implemented}}
@@ -783,16 +753,10 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc parallel loop bind(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
- // expected-warning at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +1{{expected expression}}
+#pragma acc kernels loop num_gangs()
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc parallel loop num_gangs()
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(1) num_gangs(2)
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel loop num_gangs(1) num_gangs(2)
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +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 at +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 at +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 at +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 at +1{{expected expression}}
+#pragma acc kernels loop num_workers()
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc parallel loop num_workers()
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
+#pragma acc kernels loop num_workers(1, 2)
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +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 at +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 at +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 at +1{{expected expression}}
+#pragma acc kernels loop vector_length()
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +1{{expected expression}}
+#pragma acc parallel loop vector_length()
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
+#pragma acc kernels loop vector_length(1, 2)
+ for(int i = 5; i < 10;++i);
+
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
+#pragma acc parallel loop vector_length(1, 2)
+ for(int i = 5; i < 10;++i);
+}
More information about the cfe-commits
mailing list