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
|
// 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;
// Valid cases:
#pragma acc parallel loop private(LocalInt, LocalPointer, LocalArray)
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(LocalArray)
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(LocalArray[2])
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(LocalComposite)
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(LocalComposite.EnumMember)
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(LocalComposite.ScalarMember)
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(LocalComposite.ArrayMember)
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(LocalComposite.ArrayMember[5])
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(LocalComposite.PointerMember)
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(GlobalInt, GlobalArray, GlobalPointer, GlobalComposite)
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(GlobalArray[2], GlobalPointer[2], GlobalComposite.CompositeMember.A)
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(LocalComposite, GlobalComposite)
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(IntParam, PointerParam, ArrayParam, CompositeParam) private(IntParamRef)
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(PointerParam[IntParam], ArrayParam[IntParam], CompositeParam.CompositeMember.A)
for(int i = 0; i < 5; ++i);
// 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 kernels loop private(*I)
for(int i = 0; i < 5; ++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 loop private(GlobalInt + IntParam)
for(int i = 0; i < 5; ++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 serial loop private(+GlobalInt)
for(int i = 0; i < 5; ++i);
}
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 kernels loop private(+t)
for(int i = 0; i < 5; ++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 loop private(+I)
for(int i = 0; i < 5; ++i);
// 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 serial loop private(I)
for(int i = 0; i < 5; ++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 kernels loop private(t, I)
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(arrayT)
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(TemplComp)
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(TemplComp.PointerMember[5])
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(TemplComp.PointerMember[5]) private(TemplComp)
for(int i = 0; i < 5; ++i);
int *Pointer;
#pragma acc serial loop private(Pointer[:I])
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(Pointer[:t])
for(int i = 0; i < 5; ++i);
// expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
#pragma acc parallel loop private(Pointer[1:])
for(int i = 0; i < 5; ++i);
}
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 serial loop private(I)
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(NTTP_REF)
for(int i = 0; i < 5; ++i);
}
struct S {
int ThisMember;
int ThisMemberArray[5];
void foo();
};
void S::foo() {
#pragma acc parallel loop private(ThisMember, this->ThisMemberArray[1])
for(int i = 0; i < 5; ++i);
#pragma acc serial loop private(ThisMemberArray[1:2])
for(int i = 0; i < 5; ++i);
#pragma acc kernels loop private(this)
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop private(ThisMember, this->ThisMember)
for(int i = 0; i < 5; ++i);
}
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
}
|
