Skip to content

Commit fa20b5d

Browse files
erichkeaneerichkeane
committed
[OpenACC] 'if' and 'self' clause implementation for Combined Constructs
These two are identical to how they work for compute constructs, so this patch enables them and ensures there is sufficient testing.
1 parent 4e330fa commit fa20b5d

12 files changed

+557
-42
lines changed

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -626,10 +626,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
626626

627627
OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause(
628628
SemaOpenACC::OpenACCParsedClause &Clause) {
629-
// Restrictions only properly implemented on 'compute' constructs, and
630-
// 'compute' constructs are the only construct that can do anything with
631-
// this yet, so skip/treat as unimplemented in this case.
632-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
629+
// Restrictions only properly implemented on 'compute'/'combined' constructs,
630+
// and 'compute'/'combined' constructs are the only construct that can do
631+
// anything with this yet, so skip/treat as unimplemented in this case.
632+
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
633+
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
633634
return isNotImplemented();
634635

635636
// There is no prose in the standard that says duplicates aren't allowed,
@@ -662,7 +663,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause(
662663
// Restrictions only properly implemented on 'compute' constructs, and
663664
// 'compute' constructs are the only construct that can do anything with
664665
// this yet, so skip/treat as unimplemented in this case.
665-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
666+
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
667+
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
666668
return isNotImplemented();
667669

668670
// TODO OpenACC: When we implement this for 'update', this takes a

‎clang/test/AST/ast-print-openacc-combined-construct.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -71,5 +71,19 @@ void foo() {
7171
#pragma acc kernels loop dtype(AnotherIdent)
7272
for(int i = 0;i<5;++i);
7373

74+
int i;
75+
float array[5];
76+
77+
// CHECK: #pragma acc parallel self(i == 3)
78+
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
79+
// CHECK-NEXT: ;
80+
#pragma acc parallel loop self(i == 3)
81+
for(int i = 0;i<5;++i);
82+
83+
// CHECK: #pragma acc kernels if(i == array[1])
84+
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
85+
// CHECK-NEXT: ;
86+
#pragma acc kernels if(i == array[1])
87+
for(int i = 0;i<5;++i);
7488

7589
}

‎clang/test/ParserOpenACC/parse-clauses.c

Lines changed: 3 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -266,11 +266,9 @@ void IfClause() {
266266
}
267267

268268
void SelfClause() {
269-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
270269
#pragma acc serial loop self
271270
for(int i = 0; i < 5;++i) {}
272271

273-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
274272
#pragma acc serial loop self, seq
275273
for(int i = 0; i < 5;++i) {}
276274

@@ -293,18 +291,15 @@ void SelfClause() {
293291
#pragma acc serial loop self(, seq
294292
for(int i = 0; i < 5;++i) {}
295293

296-
// expected-error@+2{{expected identifier}}
297-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
294+
// expected-error@+1{{expected identifier}}
298295
#pragma acc serial loop self)
299296
for(int i = 0; i < 5;++i) {}
300297

301-
// expected-error@+2{{expected identifier}}
302-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
298+
// expected-error@+1{{expected identifier}}
303299
#pragma acc serial loop self) seq
304300
for(int i = 0; i < 5;++i) {}
305301

306-
// expected-error@+2{{expected identifier}}
307-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented, clause ignored}}
302+
// expected-error@+1{{expected identifier}}
308303
#pragma acc serial loop self), seq
309304
for(int i = 0; i < 5;++i) {}
310305

‎clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c

Lines changed: 0 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -60,12 +60,8 @@ void uses() {
6060
// expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
6161
#pragma acc parallel loop auto default(none)
6262
for(unsigned i = 0; i < 5; ++i);
63-
// TODOexpected-error@+1{{OpenACC 'if' clause is not valid on 'parallel loop' directive}}
64-
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
6563
#pragma acc parallel loop auto if(1)
6664
for(unsigned i = 0; i < 5; ++i);
67-
// TODOexpected-error@+1{{OpenACC 'self' clause is not valid on 'parallel loop' directive}}
68-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented}}
6965
#pragma acc parallel loop auto self
7066
for(unsigned i = 0; i < 5; ++i);
7167
// TODOexpected-error@+1{{OpenACC 'copy' clause is not valid on 'parallel loop' directive}}
@@ -234,12 +230,8 @@ void uses() {
234230
// expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
235231
#pragma acc parallel loop default(none) auto
236232
for(unsigned i = 0; i < 5; ++i);
237-
// TODOexpected-error@+1{{OpenACC 'if' clause is not valid on 'parallel loop' directive}}
238-
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
239233
#pragma acc parallel loop if(1) auto
240234
for(unsigned i = 0; i < 5; ++i);
241-
// TODOexpected-error@+1{{OpenACC 'self' clause is not valid on 'parallel loop' directive}}
242-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented}}
243235
#pragma acc parallel loop self auto
244236
for(unsigned i = 0; i < 5; ++i);
245237
// TODOexpected-error@+1{{OpenACC 'copy' clause is not valid on 'parallel loop' directive}}
@@ -409,12 +401,8 @@ void uses() {
409401
// expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
410402
#pragma acc parallel loop independent default(none)
411403
for(unsigned i = 0; i < 5; ++i);
412-
// TODOexpected-error@+1{{OpenACC 'if' clause is not valid on 'parallel loop' directive}}
413-
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
414404
#pragma acc parallel loop independent if(1)
415405
for(unsigned i = 0; i < 5; ++i);
416-
// TODOexpected-error@+1{{OpenACC 'self' clause is not valid on 'parallel loop' directive}}
417-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented}}
418406
#pragma acc parallel loop independent self
419407
for(unsigned i = 0; i < 5; ++i);
420408
// TODOexpected-error@+1{{OpenACC 'copy' clause is not valid on 'parallel loop' directive}}
@@ -583,12 +571,8 @@ void uses() {
583571
// expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
584572
#pragma acc parallel loop default(none) independent
585573
for(unsigned i = 0; i < 5; ++i);
586-
// TODOexpected-error@+1{{OpenACC 'if' clause is not valid on 'parallel loop' directive}}
587-
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
588574
#pragma acc parallel loop if(1) independent
589575
for(unsigned i = 0; i < 5; ++i);
590-
// TODOexpected-error@+1{{OpenACC 'self' clause is not valid on 'parallel loop' directive}}
591-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented}}
592576
#pragma acc parallel loop self independent
593577
for(unsigned i = 0; i < 5; ++i);
594578
// TODOexpected-error@+1{{OpenACC 'copy' clause is not valid on 'parallel loop' directive}}
@@ -764,12 +748,8 @@ void uses() {
764748
// expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
765749
#pragma acc parallel loop seq default(none)
766750
for(unsigned i = 0; i < 5; ++i);
767-
// TODOexpected-error@+1{{OpenACC 'if' clause is not valid on 'parallel loop' directive}}
768-
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
769751
#pragma acc parallel loop seq if(1)
770752
for(unsigned i = 0; i < 5; ++i);
771-
// TODOexpected-error@+1{{OpenACC 'self' clause is not valid on 'parallel loop' directive}}
772-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented}}
773753
#pragma acc parallel loop seq self
774754
for(unsigned i = 0; i < 5; ++i);
775755
// TODOexpected-error@+1{{OpenACC 'copy' clause is not valid on 'parallel loop' directive}}
@@ -944,12 +924,8 @@ void uses() {
944924
// expected-warning@+1{{OpenACC clause 'default' not yet implemented}}
945925
#pragma acc parallel loop default(none) seq
946926
for(unsigned i = 0; i < 5; ++i);
947-
// TODOexpected-error@+1{{OpenACC 'if' clause is not valid on 'parallel loop' directive}}
948-
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
949927
#pragma acc parallel loop if(1) seq
950928
for(unsigned i = 0; i < 5; ++i);
951-
// TODOexpected-error@+1{{OpenACC 'self' clause is not valid on 'parallel loop' directive}}
952-
// expected-warning@+1{{OpenACC clause 'self' not yet implemented}}
953929
#pragma acc parallel loop self seq
954930
for(unsigned i = 0; i < 5; ++i);
955931
// TODOexpected-error@+1{{OpenACC 'copy' clause is not valid on 'parallel loop' directive}}

‎clang/test/SemaOpenACC/combined-construct-if-ast.cpp

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
2+
3+
// Test this with PCH.
4+
// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
5+
// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
6+
7+
#ifndef PCH_HELPER
8+
#define PCH_HELPER
9+
void NormalFunc(int j, float f) {
10+
// CHECK: FunctionDecl{{.*}}NormalFunc
11+
// CHECK-NEXT: ParmVarDecl
12+
// CHECK-NEXT: ParmVarDecl
13+
// CHECK-NEXT: CompoundStmt
14+
#pragma acc kernels loop if( j < f)
15+
for (unsigned i = 0; i < 5; ++i);
16+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop
17+
// CHECK-NEXT: if clause
18+
// CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
19+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <IntegralToFloating>
20+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
21+
// CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'j' 'int'
22+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <LValueToRValue>
23+
// CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float'
24+
// CHECK-NEXT: ForStmt
25+
// CHECK: NullStmt
26+
27+
}
28+
29+
template<typename T>
30+
void TemplFunc() {
31+
// CHECK: FunctionTemplateDecl{{.*}}TemplFunc
32+
// CHECK-NEXT: TemplateTypeParmDecl
33+
34+
// Match the prototype:
35+
// CHECK-NEXT: FunctionDecl{{.*}}TemplFunc
36+
// CHECK-NEXT: CompoundStmt
37+
38+
#pragma acc parallel loop if(T::SomeFloat < typename T::IntTy{})
39+
for (unsigned i = 0; i < 5; ++i);
40+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop
41+
// CHECK-NEXT: if clause
42+
// CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '<'
43+
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
44+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
45+
// CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename T::IntTy' 'typename T::IntTy'
46+
// CHECK-NEXT: InitListExpr{{.*}} 'void'
47+
// CHECK-NEXT: ForStmt
48+
// CHECK: NullStmt
49+
50+
#pragma acc serial loop if(typename T::IntTy{})
51+
for (unsigned i = 0; i < 5; ++i);
52+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop
53+
// CHECK-NEXT: if clause
54+
// CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename T::IntTy' 'typename T::IntTy'
55+
// CHECK-NEXT: InitListExpr{{.*}} 'void'
56+
// CHECK-NEXT: ForStmt
57+
// CHECK: NullStmt
58+
59+
#pragma acc kernels loop if(T::SomeFloat)
60+
for (unsigned i = 0; i < 5; ++i);
61+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop
62+
// CHECK-NEXT: if clause
63+
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
64+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
65+
// CHECK-NEXT: ForStmt
66+
// CHECK: NullStmt
67+
68+
#pragma acc parallel loop if(T::BC)
69+
for (unsigned i = 0; i < 5; ++i);
70+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop
71+
// CHECK-NEXT: if clause
72+
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
73+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
74+
// CHECK-NEXT: ForStmt
75+
// CHECK: NullStmt
76+
77+
// Match the instantiation:
78+
// CHECK: FunctionDecl{{.*}}TemplFunc{{.*}}implicit_instantiation
79+
// CHECK-NEXT: TemplateArgument type 'InstTy'
80+
// CHECK-NEXT: RecordType{{.*}} 'InstTy'
81+
// CHECK-NEXT: CXXRecord{{.*}} 'InstTy'
82+
// CHECK-NEXT: CompoundStmt
83+
84+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop
85+
// CHECK-NEXT: if clause
86+
// CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
87+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <LValueToRValue>
88+
// CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
89+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
90+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' <IntegralToFloating>
91+
// CHECK-NEXT: CXXFunctionalCastExpr{{.*}}'typename InstTy::IntTy':'int' functional cast to typename struct InstTy::IntTy <NoOp>
92+
// CHECK-NEXT: InitListExpr {{.*}}'typename InstTy::IntTy':'int'
93+
// CHECK-NEXT: ForStmt
94+
// CHECK: NullStmt
95+
96+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}serial loop
97+
// CHECK-NEXT: if clause
98+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <IntegralToBoolean>
99+
// CHECK-NEXT: CXXFunctionalCastExpr{{.*}}'typename InstTy::IntTy':'int' functional cast to typename struct InstTy::IntTy <NoOp>
100+
// CHECK-NEXT: InitListExpr {{.*}}'typename InstTy::IntTy':'int'
101+
// CHECK-NEXT: ForStmt
102+
// CHECK: NullStmt
103+
104+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}kernels loop
105+
// CHECK-NEXT: if clause
106+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
107+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
108+
// CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
109+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
110+
// CHECK-NEXT: ForStmt
111+
// CHECK: NullStmt
112+
113+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}}parallel loop
114+
// CHECK-NEXT: if clause
115+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'bool' <UserDefinedConversion>
116+
// CHECK-NEXT: CXXMemberCallExpr{{.*}} 'bool'
117+
// CHECK-NEXT: MemberExpr{{.*}} .operator bool
118+
// CHECK-NEXT: DeclRefExpr{{.*}} 'const BoolConversion' lvalue Var{{.*}} 'BC' 'const BoolConversion'
119+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
120+
// CHECK-NEXT: ForStmt
121+
// CHECK: NullStmt
122+
123+
}
124+
125+
struct BoolConversion{ operator bool() const;};
126+
struct InstTy {
127+
using IntTy = int;
128+
static constexpr float SomeFloat = 5.0;
129+
static constexpr BoolConversion BC;
130+
};
131+
132+
void Instantiate() {
133+
TemplFunc<InstTy>();
134+
}
135+
#endif

‎clang/test/SemaOpenACC/combined-construct-if-clause.c

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
// RUN: %clang_cc1 %s -fopenacc -verify
2+
3+
void BoolExpr(int *I, float *F) {
4+
5+
typedef struct {} SomeStruct;
6+
int Array[5];
7+
8+
struct C{};
9+
// expected-error@+1{{expected expression}}
10+
#pragma acc parallel loop if (struct C f())
11+
for (unsigned i = 0; i < 5; ++i);
12+
13+
// expected-error@+1{{unexpected type name 'SomeStruct': expected expression}}
14+
#pragma acc serial loop if (SomeStruct)
15+
for (unsigned i = 0; i < 5; ++i);
16+
17+
// expected-error@+1{{unexpected type name 'SomeStruct': expected expression}}
18+
#pragma acc serial loop if (SomeStruct())
19+
for (unsigned i = 0; i < 5; ++i);
20+
21+
SomeStruct S;
22+
// expected-error@+1{{statement requires expression of scalar type ('SomeStruct' invalid)}}
23+
#pragma acc serial loop if (S)
24+
for (unsigned i = 0; i < 5; ++i);
25+
26+
// expected-warning@+1{{address of array 'Array' will always evaluate to 'true'}}
27+
#pragma acc kernels loop if (Array)
28+
for (unsigned i = 0; i < 5; ++i);
29+
30+
// expected-warning@+4{{incompatible pointer types assigning to 'int *' from 'float *'}}
31+
// expected-warning@+3{{using the result of an assignment as a condition without parentheses}}
32+
// expected-note@+2{{place parentheses around the assignment to silence this warning}}
33+
// expected-note@+1{{use '==' to turn this assignment into an equality comparison}}
34+
#pragma acc kernels loop if (I = F)
35+
for (unsigned i = 0; i < 5; ++i);
36+
37+
#pragma acc parallel loop if (I)
38+
for (unsigned i = 0; i < 5; ++i);
39+
40+
#pragma acc serial loop if (F)
41+
for (unsigned i = 0; i < 5; ++i);
42+
43+
#pragma acc kernels loop if (*I < *F)
44+
for (unsigned i = 0; i < 5; ++i);
45+
46+
// expected-warning@+2{{OpenACC construct 'data' not yet implemented}}
47+
// expected-warning@+1{{OpenACC clause 'if' not yet implemented}}
48+
#pragma acc data if (*I < *F)
49+
for (unsigned i = 0; i < 5; ++i);
50+
#pragma acc parallel loop if (*I < *F)
51+
for(int i = 0; i < 5; ++i);
52+
#pragma acc serial loop if (*I < *F)
53+
for(int i = 0; i < 5; ++i);
54+
#pragma acc kernels loop if (*I < *F)
55+
for(int i = 0; i < 5; ++i);
56+
57+
// expected-error@+1{{OpenACC 'if' clause is not valid on 'loop' directive}}
58+
#pragma acc loop if(I)
59+
for(int i = 5; i < 10;++i);
60+
}

‎clang/test/SemaOpenACC/combined-construct-if-clause.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// RUN: %clang_cc1 %s -fopenacc -verify
2+
3+
struct NoBoolConversion{};
4+
struct BoolConversion{
5+
operator bool();
6+
};
7+
8+
template <typename T, typename U>
9+
void BoolExpr() {
10+
11+
// expected-error@+1{{value of type 'NoBoolConversion' is not contextually convertible to 'bool'}}
12+
#pragma acc parallel loop if (NoBoolConversion{})
13+
for (unsigned i = 0; i < 5; ++i);
14+
15+
// expected-error@+2{{no member named 'NotValid' in 'NoBoolConversion'}}
16+
// expected-note@#INST{{in instantiation of function template specialization}}
17+
#pragma acc serial loop if (T::NotValid)
18+
for (unsigned i = 0; i < 5; ++i);
19+
20+
#pragma acc kernels loop if (BoolConversion{})
21+
for (unsigned i = 0; i < 5; ++i);
22+
23+
// expected-error@+1{{value of type 'NoBoolConversion' is not contextually convertible to 'bool'}}
24+
#pragma acc serial loop if (T{})
25+
for (unsigned i = 0; i < 5; ++i);
26+
27+
#pragma acc parallel loop if (U{})
28+
for (unsigned i = 0; i < 5; ++i);
29+
}
30+
31+
void Instantiate() {
32+
BoolExpr<NoBoolConversion, BoolConversion>(); // #INST
33+
}

0 commit comments

Comments
 (0)