Skip to content

Commit eb257fe

Browse files
erichkeaneerichkeane
committed
[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.
1 parent 95a4d30 commit eb257fe

10 files changed

+587
-59
lines changed

‎clang/lib/Sema/SemaOpenACC.cpp‎

Lines changed: 1 addition & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -694,14 +694,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause(
694694

695695
OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
696696
SemaOpenACC::OpenACCParsedClause &Clause) {
697-
// Restrictions only properly implemented on 'compute' constructs, and
698-
// 'compute' constructs are the only construct that can do anything with
699-
// this yet, so skip/treat as unimplemented in this case.
700-
// TODO OpenACC: Remove this check when we have combined constructs for this
701-
// clause.
702-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
703-
return isNotImplemented();
704-
705697
// There is no prose in the standard that says duplicates aren't allowed,
706698
// but this diagnostic is present in other compilers, as well as makes
707699
// sense.
@@ -730,6 +722,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
730722
// OpenACC 3.3 Section 2.5.4:
731723
// A reduction clause may not appear on a parallel construct with a
732724
// num_gangs clause that has more than one argument.
725+
// TODO: OpenACC: Reduction on Combined Construct needs to do this too.
733726
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel &&
734727
Clause.getIntExprs().size() > 1) {
735728
auto *Parallel =
@@ -751,13 +744,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
751744

752745
OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
753746
SemaOpenACC::OpenACCParsedClause &Clause) {
754-
// Restrictions only properly implemented on 'compute' constructs, and
755-
// 'compute' constructs are the only construct that can do anything with
756-
// this yet, so skip/treat as unimplemented in this case.
757-
// TODO: OpenACC: Remove when we get combined constructs.
758-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
759-
return isNotImplemented();
760-
761747
// There is no prose in the standard that says duplicates aren't allowed,
762748
// but this diagnostic is present in other compilers, as well as makes
763749
// sense.
@@ -773,13 +759,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
773759

774760
OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
775761
SemaOpenACC::OpenACCParsedClause &Clause) {
776-
// Restrictions only properly implemented on 'compute' constructs, and
777-
// 'compute' constructs are the only construct that can do anything with
778-
// this yet, so skip/treat as unimplemented in this case.
779-
// TODO: OpenACC: Remove when we get combined constructs.
780-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
781-
return isNotImplemented();
782-
783762
// There is no prose in the standard that says duplicates aren't allowed,
784763
// but this diagnostic is present in other compilers, as well as makes
785764
// sense.

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

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -224,4 +224,22 @@ void foo() {
224224
for(int i = 0;i<5;++i)
225225
for(int i = 0;i<5;++i);
226226

227+
// CHECK: #pragma acc parallel loop num_gangs(i, (int)array[2])
228+
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
229+
// CHECK-NEXT: ;
230+
#pragma acc parallel loop num_gangs(i, (int)array[2])
231+
for(int i = 0;i<5;++i);
232+
233+
// CHECK: #pragma acc parallel loop num_workers(i)
234+
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
235+
// CHECK-NEXT: ;
236+
#pragma acc parallel loop num_workers(i)
237+
for(int i = 0;i<5;++i);
238+
239+
// CHECK: #pragma acc parallel loop vector_length((int)array[1])
240+
// CHECK-NEXT: for (int i = 0; i < 5; ++i)
241+
// CHECK-NEXT: ;
242+
#pragma acc parallel loop vector_length((int)array[1])
243+
for(int i = 0;i<5;++i);
244+
227245
}

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

Lines changed: 0 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -134,16 +134,10 @@ void uses() {
134134
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
135135
#pragma acc parallel loop auto bind(Var)
136136
for(unsigned i = 0; i < 5; ++i);
137-
// TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
138-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
139137
#pragma acc parallel loop auto vector_length(1)
140138
for(unsigned i = 0; i < 5; ++i);
141-
// TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
142-
// expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
143139
#pragma acc parallel loop auto num_gangs(1)
144140
for(unsigned i = 0; i < 5; ++i);
145-
// TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
146-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
147141
#pragma acc parallel loop auto num_workers(1)
148142
for(unsigned i = 0; i < 5; ++i);
149143
// expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -261,16 +255,10 @@ void uses() {
261255
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
262256
#pragma acc parallel loop bind(Var) auto
263257
for(unsigned i = 0; i < 5; ++i);
264-
// TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
265-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
266258
#pragma acc parallel loop vector_length(1) auto
267259
for(unsigned i = 0; i < 5; ++i);
268-
// TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
269-
// expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
270260
#pragma acc parallel loop num_gangs(1) auto
271261
for(unsigned i = 0; i < 5; ++i);
272-
// TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
273-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
274262
#pragma acc parallel loop num_workers(1) auto
275263
for(unsigned i = 0; i < 5; ++i);
276264
// expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -389,16 +377,10 @@ void uses() {
389377
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
390378
#pragma acc parallel loop independent bind(Var)
391379
for(unsigned i = 0; i < 5; ++i);
392-
// TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
393-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
394380
#pragma acc parallel loop independent vector_length(1)
395381
for(unsigned i = 0; i < 5; ++i);
396-
// TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
397-
// expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
398382
#pragma acc parallel loop independent num_gangs(1)
399383
for(unsigned i = 0; i < 5; ++i);
400-
// TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
401-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
402384
#pragma acc parallel loop independent num_workers(1)
403385
for(unsigned i = 0; i < 5; ++i);
404386
// expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -516,16 +498,10 @@ void uses() {
516498
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
517499
#pragma acc parallel loop bind(Var) independent
518500
for(unsigned i = 0; i < 5; ++i);
519-
// TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
520-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
521501
#pragma acc parallel loop vector_length(1) independent
522502
for(unsigned i = 0; i < 5; ++i);
523-
// TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
524-
// expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
525503
#pragma acc parallel loop num_gangs(1) independent
526504
for(unsigned i = 0; i < 5; ++i);
527-
// TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
528-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
529505
#pragma acc parallel loop num_workers(1) independent
530506
for(unsigned i = 0; i < 5; ++i);
531507
// expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -650,16 +626,10 @@ void uses() {
650626
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
651627
#pragma acc parallel loop seq bind(Var)
652628
for(unsigned i = 0; i < 5; ++i);
653-
// TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
654-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
655629
#pragma acc parallel loop seq vector_length(1)
656630
for(unsigned i = 0; i < 5; ++i);
657-
// TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
658-
// expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
659631
#pragma acc parallel loop seq num_gangs(1)
660632
for(unsigned i = 0; i < 5; ++i);
661-
// TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
662-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
663633
#pragma acc parallel loop seq num_workers(1)
664634
for(unsigned i = 0; i < 5; ++i);
665635
// expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -783,16 +753,10 @@ void uses() {
783753
// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
784754
#pragma acc parallel loop bind(Var) seq
785755
for(unsigned i = 0; i < 5; ++i);
786-
// TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
787-
// expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
788756
#pragma acc parallel loop vector_length(1) seq
789757
for(unsigned i = 0; i < 5; ++i);
790-
// TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
791-
// expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
792758
#pragma acc parallel loop num_gangs(1) seq
793759
for(unsigned i = 0; i < 5; ++i);
794-
// TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
795-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
796760
#pragma acc parallel loop num_workers(1) seq
797761
for(unsigned i = 0; i < 5; ++i);
798762
// expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -195,7 +195,6 @@ void uses() {
195195
// expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial loop' directive}}
196196
#pragma acc serial loop device_type(*) num_gangs(1)
197197
for(int i = 0; i < 5; ++i);
198-
// expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented, clause ignored}}
199198
#pragma acc parallel loop device_type(*) num_workers(1)
200199
for(int i = 0; i < 5; ++i);
201200
// expected-error@+2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a 'serial loop' construct}}

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

Lines changed: 121 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,121 @@
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+
int some_int();
10+
short some_short();
11+
long some_long();
12+
struct CorrectConvert {
13+
operator int();
14+
} Convert;
15+
16+
17+
void NormalUses() {
18+
// CHECK: FunctionDecl{{.*}}NormalUses
19+
// CHECK-NEXT: CompoundStmt
20+
21+
#pragma acc parallel loop num_gangs(some_int(), some_long(), some_short())
22+
for (unsigned i = 0; i < 5; ++i);
23+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
24+
// CHECK-NEXT: num_gangs clause
25+
// CHECK-NEXT: CallExpr{{.*}}'int'
26+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
27+
// CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
28+
// CHECK-NEXT: CallExpr{{.*}}'long'
29+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay>
30+
// CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()'
31+
// CHECK-NEXT: CallExpr{{.*}}'short'
32+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
33+
// CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
34+
// CHECK-NEXT: ForStmt
35+
// CHECK: NullStmt
36+
37+
#pragma acc kernels loop num_gangs(some_int())
38+
for (unsigned i = 0; i < 5; ++i);
39+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
40+
// CHECK-NEXT: num_gangs clause
41+
// CHECK-NEXT: CallExpr{{.*}}'int'
42+
// CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
43+
// CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
44+
// CHECK-NEXT: ForStmt
45+
// CHECK: NullStmt
46+
}
47+
48+
template<typename T, typename U>
49+
void TemplUses(T t, U u) {
50+
// CHECK-NEXT: FunctionTemplateDecl
51+
// CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
52+
// CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U
53+
// CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)'
54+
// CHECK-NEXT: ParmVarDecl{{.*}} t 'T'
55+
// CHECK-NEXT: ParmVarDecl{{.*}} u 'U'
56+
// CHECK-NEXT: CompoundStmt
57+
58+
#pragma acc kernels loop num_gangs(u)
59+
for (unsigned i = 0; i < 5; ++i);
60+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
61+
// CHECK-NEXT: num_gangs clause
62+
// CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
63+
// CHECK-NEXT: ForStmt
64+
// CHECK: NullStmt
65+
66+
#pragma acc parallel loop num_gangs(u, U::value)
67+
for (unsigned i = 0; i < 5; ++i);
68+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
69+
// CHECK-NEXT: num_gangs clause
70+
// CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
71+
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
72+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
73+
// CHECK-NEXT: ForStmt
74+
// CHECK: NullStmt
75+
76+
// Check the instantiated versions of the above.
77+
// CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
78+
// CHECK-NEXT: TemplateArgument type 'CorrectConvert'
79+
// CHECK-NEXT: RecordType{{.*}} 'CorrectConvert'
80+
// CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert'
81+
// CHECK-NEXT: TemplateArgument type 'HasInt'
82+
// CHECK-NEXT: RecordType{{.*}} 'HasInt'
83+
// CHECK-NEXT: CXXRecord{{.*}} 'HasInt'
84+
// CHECK-NEXT: ParmVarDecl{{.*}} t 'CorrectConvert'
85+
// CHECK-NEXT: ParmVarDecl{{.*}} u 'HasInt'
86+
// CHECK-NEXT: CompoundStmt
87+
88+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
89+
// CHECK-NEXT: num_gangs clause
90+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
91+
// CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
92+
// CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
93+
// CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
94+
// CHECK-NEXT: ForStmt
95+
// CHECK: NullStmt
96+
97+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
98+
// CHECK-NEXT: num_gangs clause
99+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
100+
// CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
101+
// CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
102+
// CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
103+
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
104+
// CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
105+
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
106+
// CHECK-NEXT: ForStmt
107+
// CHECK: NullStmt
108+
}
109+
110+
struct HasInt {
111+
using IntTy = int;
112+
using ShortTy = short;
113+
static constexpr int value = 1;
114+
115+
operator char();
116+
};
117+
118+
void Inst() {
119+
TemplUses<CorrectConvert, HasInt>({}, {});
120+
}
121+
#endif // PCH_HELPER

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
// RUN: %clang_cc1 %s -fopenacc -verify
2+
3+
short getS();
4+
float getF();
5+
void Test() {
6+
#pragma acc kernels loop num_gangs(1)
7+
for(int i = 5; i < 10;++i);
8+
9+
// expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial loop' directive}}
10+
#pragma acc serial loop num_gangs(1)
11+
for(int i = 5; i < 10;++i);
12+
13+
#pragma acc parallel loop num_gangs(1)
14+
for(int i = 5; i < 10;++i);
15+
16+
// expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type}}
17+
#pragma acc parallel loop num_gangs(getF())
18+
for(int i = 5; i < 10;++i);
19+
20+
// expected-error@+1{{expected expression}}
21+
#pragma acc kernels loop num_gangs()
22+
for(int i = 5; i < 10;++i);
23+
24+
// expected-error@+1{{expected expression}}
25+
#pragma acc parallel loop num_gangs()
26+
for(int i = 5; i < 10;++i);
27+
28+
// expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}}
29+
// expected-note@+1{{previous clause is here}}
30+
#pragma acc kernels loop num_gangs(1) num_gangs(2)
31+
for(int i = 5; i < 10;++i);
32+
33+
// expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}}
34+
// expected-note@+1{{previous clause is here}}
35+
#pragma acc parallel loop num_gangs(1) num_gangs(2)
36+
for(int i = 5; i < 10;++i);
37+
38+
// expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels loop' directive expects maximum of 1, 2 were provided}}
39+
#pragma acc kernels loop num_gangs(1, getS())
40+
for(int i = 5; i < 10;++i);
41+
42+
// expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel loop' directive expects maximum of 3, 4 were provided}}
43+
#pragma acc parallel loop num_gangs(getS(), 1, getS(), 1)
44+
for(int i = 5; i < 10;++i);
45+
}

0 commit comments

Comments
 (0)