Skip to content

Commit 8d6c73c

Browse files
erichkeaneerichkeane
committed
[OpenACC] enable 'deviceptr' for combined constructs.
This is another clause whose implementation is identical for combined constructs as with compute constructs, so this adds tests and enables it.
1 parent 43ee6f7 commit 8d6c73c

6 files changed

+264
-16
lines changed

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -961,10 +961,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
961961

962962
OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
963963
SemaOpenACC::OpenACCParsedClause &Clause) {
964-
// Restrictions only properly implemented on 'compute' constructs, and
965-
// 'compute' constructs are the only construct that can do anything with
966-
// this yet, so skip/treat as unimplemented in this case.
967-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
964+
// Restrictions only properly implemented on 'compute'/'combined' constructs,
965+
// and 'compute'/'combined' constructs are the only construct that can do
966+
// anything with this yet, so skip/treat as unimplemented in this case.
967+
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
968+
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
968969
return isNotImplemented();
969970

970971
// ActOnVar ensured that everything is a valid variable reference, but we

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -131,4 +131,10 @@ void foo() {
131131
// CHECK: #pragma acc kernels loop present(i, array[1], array, array[1:2])
132132
#pragma acc kernels loop present(i, array[1], array, array[1:2])
133133
for(int i = 0;i<5;++i);
134+
135+
float *arrayPtr[5];
136+
137+
// CHECK: #pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])
138+
#pragma acc kernels loop deviceptr(iPtr, arrayPtr[0])
139+
for(int i = 0;i<5;++i);
134140
}

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

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -92,8 +92,6 @@ void uses() {
9292
// expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
9393
#pragma acc parallel loop auto device(VarPtr)
9494
for(unsigned i = 0; i < 5; ++i);
95-
// TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
96-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
9795
#pragma acc parallel loop auto deviceptr(VarPtr)
9896
for(unsigned i = 0; i < 5; ++i);
9997
// expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -253,8 +251,6 @@ void uses() {
253251
// expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
254252
#pragma acc parallel loop device(VarPtr) auto
255253
for(unsigned i = 0; i < 5; ++i);
256-
// TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
257-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
258254
#pragma acc parallel loop deviceptr(VarPtr) auto
259255
for(unsigned i = 0; i < 5; ++i);
260256
// expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -415,8 +411,6 @@ void uses() {
415411
// expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
416412
#pragma acc parallel loop independent device(VarPtr)
417413
for(unsigned i = 0; i < 5; ++i);
418-
// TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
419-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
420414
#pragma acc parallel loop independent deviceptr(VarPtr)
421415
for(unsigned i = 0; i < 5; ++i);
422416
// expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -576,8 +570,6 @@ void uses() {
576570
// expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
577571
#pragma acc parallel loop device(VarPtr) independent
578572
for(unsigned i = 0; i < 5; ++i);
579-
// TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
580-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
581573
#pragma acc parallel loop deviceptr(VarPtr) independent
582574
for(unsigned i = 0; i < 5; ++i);
583575
// expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -744,8 +736,6 @@ void uses() {
744736
// expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
745737
#pragma acc parallel loop seq device(VarPtr)
746738
for(unsigned i = 0; i < 5; ++i);
747-
// TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
748-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
749739
#pragma acc parallel loop seq deviceptr(VarPtr)
750740
for(unsigned i = 0; i < 5; ++i);
751741
// expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
@@ -911,8 +901,6 @@ void uses() {
911901
// expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
912902
#pragma acc parallel loop device(VarPtr) seq
913903
for(unsigned i = 0; i < 5; ++i);
914-
// TODOexpected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'parallel loop' directive}}
915-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
916904
#pragma acc parallel loop deviceptr(VarPtr) seq
917905
for(unsigned i = 0; i < 5; ++i);
918906
// expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}

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

Lines changed: 64 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,64 @@
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+
10+
int Global;
11+
short GlobalArray[5];
12+
13+
void NormalUses(float *PointerParam) {
14+
// CHECK: FunctionDecl{{.*}}NormalUses
15+
// CHECK: ParmVarDecl
16+
// CHECK-NEXT: CompoundStmt
17+
18+
#pragma acc parallel loop deviceptr(PointerParam)
19+
for(int i = 0; i < 5; ++i);
20+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
21+
// CHECK-NEXT: deviceptr clause
22+
// CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
23+
// CHECK-NEXT: ForStmt
24+
// CHECK: NullStmt
25+
}
26+
27+
template<typename T>
28+
void TemplUses(T *t) {
29+
// CHECK-NEXT: FunctionTemplateDecl
30+
// CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
31+
// CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T *)'
32+
// CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T *'
33+
// CHECK-NEXT: CompoundStmt
34+
35+
#pragma acc parallel loop deviceptr(t)
36+
for(int i = 0; i < 5; ++i);
37+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
38+
// CHECK-NEXT: deviceptr clause
39+
// CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 't' 'T *'
40+
// CHECK-NEXT: ForStmt
41+
// CHECK: NullStmt
42+
43+
44+
// Check the instantiated versions of the above.
45+
// CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (int *)' implicit_instantiation
46+
// CHECK-NEXT: TemplateArgument type 'int'
47+
// CHECK-NEXT: BuiltinType{{.*}} 'int'
48+
// CHECK-NEXT: ParmVarDecl{{.*}} used t 'int *'
49+
// CHECK-NEXT: CompoundStmt
50+
51+
// #pragma acc parallel loop deviceptr(t)
52+
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
53+
// CHECK-NEXT: deviceptr clause
54+
// CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 't' 'int *'
55+
// CHECK-NEXT: ForStmt
56+
// CHECK: NullStmt
57+
58+
}
59+
60+
void Inst() {
61+
int i;
62+
TemplUses(&i);
63+
}
64+
#endif

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

Lines changed: 69 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,69 @@
1+
// RUN: %clang_cc1 %s -fopenacc -verify
2+
3+
struct S {
4+
int IntMem;
5+
int *PtrMem;
6+
};
7+
8+
void uses() {
9+
int LocalInt;
10+
int *LocalPtr;
11+
int Array[5];
12+
int *PtrArray[5];
13+
struct S s;
14+
15+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
16+
#pragma acc parallel loop deviceptr(LocalInt)
17+
for (int i = 0; i < 5; ++i);
18+
19+
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
20+
#pragma acc parallel loop deviceptr(&LocalInt)
21+
for (int i = 0; i < 5; ++i);
22+
23+
#pragma acc serial loop deviceptr(LocalPtr)
24+
for (int i = 0; i < 5; ++i);
25+
26+
#pragma acc kernels loop deviceptr(LocalPtr)
27+
for (int i = 0; i < 5; ++i);
28+
29+
30+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
31+
#pragma acc kernels loop deviceptr(Array)
32+
for (int i = 0; i < 5; ++i);
33+
34+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
35+
#pragma acc parallel loop deviceptr(Array[0])
36+
for (int i = 0; i < 5; ++i);
37+
38+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
39+
// expected-note@+1{{expected variable of pointer type}}
40+
#pragma acc parallel loop deviceptr(Array[0:1])
41+
for (int i = 0; i < 5; ++i);
42+
43+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
44+
#pragma acc parallel loop deviceptr(PtrArray)
45+
for (int i = 0; i < 5; ++i);
46+
47+
#pragma acc parallel loop deviceptr(PtrArray[0])
48+
for (int i = 0; i < 5; ++i);
49+
50+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
51+
// expected-note@+1{{expected variable of pointer type}}
52+
#pragma acc parallel loop deviceptr(PtrArray[0:1])
53+
for (int i = 0; i < 5; ++i);
54+
55+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
56+
#pragma acc parallel loop deviceptr(s)
57+
for (int i = 0; i < 5; ++i);
58+
59+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
60+
#pragma acc parallel loop deviceptr(s.IntMem)
61+
for (int i = 0; i < 5; ++i);
62+
63+
#pragma acc parallel loop deviceptr(s.PtrMem)
64+
for (int i = 0; i < 5; ++i);
65+
66+
// expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
67+
#pragma acc loop deviceptr(LocalInt)
68+
for(int i = 5; i < 10;++i);
69+
}

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

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,120 @@
1+
// RUN: %clang_cc1 %s -fopenacc -verify
2+
3+
struct S {
4+
int IntMem;
5+
int *PtrMem;
6+
operator int*();
7+
};
8+
9+
void uses() {
10+
int LocalInt;
11+
int *LocalPtr;
12+
int Array[5];
13+
int *PtrArray[5];
14+
struct S s;
15+
16+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
17+
#pragma acc parallel loop deviceptr(LocalInt)
18+
for (int i = 0; i < 5; ++i);
19+
20+
#pragma acc parallel loop deviceptr(LocalPtr)
21+
for (int i = 0; i < 5; ++i);
22+
23+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
24+
#pragma acc parallel loop deviceptr(Array)
25+
for (int i = 0; i < 5; ++i);
26+
27+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
28+
#pragma acc parallel loop deviceptr(Array[0])
29+
for (int i = 0; i < 5; ++i);
30+
31+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
32+
// expected-note@+1{{expected variable of pointer type}}
33+
#pragma acc parallel loop deviceptr(Array[0:1])
34+
for (int i = 0; i < 5; ++i);
35+
36+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
37+
#pragma acc parallel loop deviceptr(PtrArray)
38+
for (int i = 0; i < 5; ++i);
39+
40+
#pragma acc parallel loop deviceptr(PtrArray[0])
41+
for (int i = 0; i < 5; ++i);
42+
43+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
44+
// expected-note@+1{{expected variable of pointer type}}
45+
#pragma acc parallel loop deviceptr(PtrArray[0:1])
46+
for (int i = 0; i < 5; ++i);
47+
48+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
49+
#pragma acc parallel loop deviceptr(s)
50+
for (int i = 0; i < 5; ++i);
51+
52+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
53+
#pragma acc parallel loop deviceptr(s.IntMem)
54+
for (int i = 0; i < 5; ++i);
55+
56+
#pragma acc parallel loop deviceptr(s.PtrMem)
57+
for (int i = 0; i < 5; ++i);
58+
}
59+
60+
template<typename T, typename TPtr, typename TStruct, auto &R1>
61+
void Templ() {
62+
T SomeInt;
63+
TPtr SomePtr;
64+
T SomeIntArray[5];
65+
TPtr SomeIntPtrArray[5];
66+
TStruct SomeStruct;
67+
68+
// expected-error@+2{{expected pointer in 'deviceptr' clause, type is 'int'}}
69+
// expected-note@#INST{{in instantiation of function template specialization}}
70+
#pragma acc parallel loop deviceptr(SomeInt)
71+
for (int i = 0; i < 5; ++i);
72+
73+
#pragma acc parallel loop deviceptr(SomePtr)
74+
for (int i = 0; i < 5; ++i);
75+
76+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
77+
#pragma acc parallel loop deviceptr(SomeIntArray)
78+
for (int i = 0; i < 5; ++i);
79+
80+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
81+
#pragma acc parallel loop deviceptr(SomeIntArray[0])
82+
for (int i = 0; i < 5; ++i);
83+
84+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
85+
// expected-note@+1{{expected variable of pointer type}}
86+
#pragma acc parallel loop deviceptr(SomeIntArray[0:1])
87+
for (int i = 0; i < 5; ++i);
88+
89+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
90+
#pragma acc parallel loop deviceptr(SomeIntPtrArray)
91+
for (int i = 0; i < 5; ++i);
92+
93+
#pragma acc parallel loop deviceptr(SomeIntPtrArray[0])
94+
for (int i = 0; i < 5; ++i);
95+
96+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
97+
// expected-note@+1{{expected variable of pointer type}}
98+
#pragma acc parallel loop deviceptr(SomeIntPtrArray[0:1])
99+
for (int i = 0; i < 5; ++i);
100+
101+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'S'}}
102+
#pragma acc parallel loop deviceptr(SomeStruct)
103+
for (int i = 0; i < 5; ++i);
104+
105+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
106+
#pragma acc parallel loop deviceptr(SomeStruct.IntMem)
107+
for (int i = 0; i < 5; ++i);
108+
109+
#pragma acc parallel loop deviceptr(SomeStruct.PtrMem)
110+
for (int i = 0; i < 5; ++i);
111+
112+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
113+
#pragma acc parallel loop deviceptr(R1)
114+
for (int i = 0; i < 5; ++i);
115+
}
116+
117+
void inst() {
118+
static constexpr int CEVar = 1;
119+
Templ<int, int*, S, CEVar>(); // #INST
120+
}

0 commit comments

Comments
 (0)