1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
|
// RUN: %clang_cc1 %s -fopenacc -verify
struct Incomplete;
enum SomeE{};
typedef struct IsComplete {
struct S { int A; } CompositeMember;
int ScalarMember;
float ArrayMember[5];
SomeE EnumMember;
char *PointerMember;
} Complete;
int GlobalInt;
float GlobalArray[5];
char *GlobalPointer;
Complete GlobalComposite;
void uses(int IntParam, char *PointerParam, float ArrayParam[5], Complete CompositeParam, int &IntParamRef) {
int LocalInt;
char *LocalPointer;
float LocalArray[5];
Complete LocalComposite;
// Check Appertainment:
#pragma acc parallel private(LocalInt)
while(true);
#pragma acc serial private(LocalInt)
while(true);
// expected-error@+1{{OpenACC 'private' clause is not valid on 'kernels' directive}}
#pragma acc kernels private(LocalInt)
while(true);
// Valid cases:
#pragma acc parallel private(LocalInt, LocalPointer, LocalArray)
while(true);
#pragma acc parallel private(LocalArray)
while(true);
#pragma acc parallel private(LocalArray[2])
while(true);
#pragma acc parallel private(LocalComposite)
while(true);
#pragma acc parallel private(LocalComposite.EnumMember)
while(true);
#pragma acc parallel private(LocalComposite.ScalarMember)
while(true);
#pragma acc parallel private(LocalComposite.ArrayMember)
while(true);
#pragma acc parallel private(LocalComposite.ArrayMember[5])
while(true);
#pragma acc parallel private(LocalComposite.PointerMember)
while(true);
#pragma acc parallel private(GlobalInt, GlobalArray, GlobalPointer, GlobalComposite)
while(true);
#pragma acc parallel private(GlobalArray[2], GlobalPointer[2], GlobalComposite.CompositeMember.A)
while(true);
#pragma acc parallel private(LocalComposite, GlobalComposite)
while(true);
#pragma acc parallel private(IntParam, PointerParam, ArrayParam, CompositeParam) private(IntParamRef)
while(true);
#pragma acc parallel private(PointerParam[IntParam], ArrayParam[IntParam], CompositeParam.CompositeMember.A)
while(true);
// Invalid cases, arbitrary expressions.
Incomplete *I;
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
#pragma acc parallel private(*I)
while(true);
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
#pragma acc parallel private(GlobalInt + IntParam)
while(true);
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
#pragma acc parallel private(+GlobalInt)
while(true);
}
template<typename T, unsigned I, typename V>
void TemplUses(T t, T (&arrayT)[I], V TemplComp) {
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
#pragma acc parallel private(+t)
while(true);
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
#pragma acc parallel private(+I)
while(true);
// NTTP's are only valid if it is a reference to something.
// expected-error@+2{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
// expected-note@#TEMPL_USES_INST{{in instantiation of}}
#pragma acc parallel private(I)
while(true);
// expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
#pragma acc parallel private(t, I)
while(true);
#pragma acc parallel private(arrayT)
while(true);
#pragma acc parallel private(TemplComp)
while(true);
#pragma acc parallel private(TemplComp.PointerMember[5])
while(true);
#pragma acc parallel private(TemplComp.PointerMember[5]) private(TemplComp)
while(true);
int *Pointer;
#pragma acc parallel private(Pointer[:I])
while(true);
#pragma acc parallel private(Pointer[:t])
while(true);
// expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
#pragma acc parallel private(Pointer[1:])
while(true);
}
template<unsigned I, auto &NTTP_REF>
void NTTP() {
// NTTP's are only valid if it is a reference to something.
// expected-error@+2{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
// expected-note@#NTTP_INST{{in instantiation of}}
#pragma acc parallel private(I)
while(true);
#pragma acc parallel private(NTTP_REF)
while(true);
}
struct S {
int ThisMember;
int ThisMemberArray[5];
void foo();
};
void S::foo() {
#pragma acc parallel private(ThisMember, this->ThisMemberArray[1])
while(true);
#pragma acc parallel private(ThisMemberArray[1:2])
while(true);
#pragma acc parallel private(this)
while(true);
#pragma acc parallel private(ThisMember, this->ThisMember)
while(true);
}
void Inst() {
static constexpr int NTTP_REFed = 1;
int i;
int Arr[5];
Complete C;
TemplUses(i, Arr, C); // #TEMPL_USES_INST
NTTP<5, NTTP_REFed>(); // #NTTP_INST
}
|
