Skip to content

Commit 754499c

Browse files
erichkeaneerichkeane
committed
[OpenACC] Enable 'deviceptr' clause sema on data construct
Another simple implementation, as it uses the same work as the previous implementation, just enabling it for this construct.
1 parent 5225f1b commit 754499c

File tree

5 files changed

+151
-7
lines changed

5 files changed

+151
-7
lines changed

‎clang/lib/Sema/SemaOpenACC.cpp‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1035,11 +1035,13 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
10351035

10361036
OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
10371037
SemaOpenACC::OpenACCParsedClause &Clause) {
1038-
// Restrictions only properly implemented on 'compute'/'combined' constructs,
1039-
// and 'compute'/'combined' constructs are the only construct that can do
1040-
// anything with this yet, so skip/treat as unimplemented in this case.
1038+
// Restrictions only properly implemented on 'compute'/'combined'/'data'
1039+
// constructs, and 'compute'/'combined'/'data' constructs are the only
1040+
// construct that can do anything with this yet, so skip/treat as
1041+
// unimplemented in this case.
10411042
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
1042-
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
1043+
!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) &&
1044+
!isOpenACCDataDirectiveKind(Clause.getDirectiveKind()))
10431045
return isNotImplemented();
10441046

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

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,4 +101,10 @@ void foo() {
101101

102102
// CHECK: #pragma acc enter data create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
103103
#pragma acc enter data create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
104+
105+
float *arrayPtr[5];
106+
107+
// CHECK: #pragma acc data default(none) deviceptr(iPtr, arrayPtr[0])
108+
#pragma acc data default(none) deviceptr(iPtr, arrayPtr[0])
109+
;
104110
}

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

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

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

Lines changed: 70 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,70 @@
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 data default(none) 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 data default(none) deviceptr(&LocalInt)
21+
for (int i = 0; i < 5; ++i);
22+
23+
24+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
25+
#pragma acc kernels loop deviceptr(Array)
26+
for (int i = 0; i < 5; ++i);
27+
28+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
29+
#pragma acc data default(none) deviceptr(Array[0])
30+
for (int i = 0; i < 5; ++i);
31+
32+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
33+
// expected-note@+1{{expected variable of pointer type}}
34+
#pragma acc data default(none) deviceptr(Array[0:1])
35+
for (int i = 0; i < 5; ++i);
36+
37+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
38+
#pragma acc data default(none) deviceptr(PtrArray)
39+
for (int i = 0; i < 5; ++i);
40+
41+
#pragma acc data default(none) deviceptr(PtrArray[0])
42+
for (int i = 0; i < 5; ++i);
43+
44+
// expected-error@+2{{OpenACC sub-array is not allowed here}}
45+
// expected-note@+1{{expected variable of pointer type}}
46+
#pragma acc data default(none) deviceptr(PtrArray[0:1])
47+
for (int i = 0; i < 5; ++i);
48+
49+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
50+
#pragma acc data default(none) deviceptr(s)
51+
for (int i = 0; i < 5; ++i);
52+
53+
// expected-error@+1{{expected pointer in 'deviceptr' clause, type is 'int'}}
54+
#pragma acc data default(none) deviceptr(s.IntMem)
55+
for (int i = 0; i < 5; ++i);
56+
57+
#pragma acc data default(none) deviceptr(s.PtrMem)
58+
for (int i = 0; i < 5; ++i);
59+
60+
// expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'enter data' directive}}
61+
#pragma acc enter data copyin(LocalInt) deviceptr(LocalInt)
62+
for(int i = 5; i < 10;++i);
63+
// expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'exit data' directive}}
64+
#pragma acc exit data copyout(LocalInt) deviceptr(LocalInt)
65+
for(int i = 5; i < 10;++i);
66+
// expected-warning@+2{{OpenACC clause 'use_device' not yet implemented}}
67+
// expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'host_data' directive}}
68+
#pragma acc host_data use_device(LocalInt) deviceptr(LocalInt)
69+
for(int i = 5; i < 10;++i);
70+
}

‎clang/test/SemaOpenACC/data-construct.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,7 @@ void HasStmt() {
2020

2121
void AtLeastOneOf() {
2222
int Var;
23+
int *VarPtr = &Var;
2324
// Data
2425
#pragma acc data copy(Var)
2526
;
@@ -33,11 +34,10 @@ void AtLeastOneOf() {
3334
;
3435
#pragma acc data present(Var)
3536
;
36-
// expected-warning@+1{{OpenACC clause 'deviceptr' not yet implemented}}
37-
#pragma acc data deviceptr(Var)
37+
#pragma acc data deviceptr(VarPtr)
3838
;
3939
// expected-warning@+1{{OpenACC clause 'attach' not yet implemented}}
40-
#pragma acc data attach(Var)
40+
#pragma acc data attach(VarPtr)
4141
;
4242
#pragma acc data default(none)
4343
;

0 commit comments

Comments
 (0)