1 // RUN: %clang_cc1 -verify -fopenmp -std=c++11 -ast-print %s | FileCheck %s 2 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s 3 // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s 4 5 // RUN: %clang_cc1 -verify -fopenmp-simd -std=c++11 -ast-print %s | FileCheck %s 6 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s 7 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s 8 // expected-no-diagnostics 9 10 #ifndef HEADER 11 #define HEADER 12 struct ST { 13 int *a; 14 }; 15 typedef int arr[10]; 16 typedef ST STarr[10]; 17 struct SA { 18 const int da[5] = { 0 }; 19 ST g[10]; 20 STarr &rg = g; 21 int i; 22 int &j = i; 23 int *k = &j; 24 int *&z = k; 25 int aa[10]; 26 arr &raa = aa; funcSA27 void func(int arg) { 28 #pragma omp target teams is_device_ptr(k) 29 {} 30 #pragma omp target teams is_device_ptr(z) 31 {} 32 #pragma omp target teams is_device_ptr(aa) // OK 33 {} 34 #pragma omp target teams is_device_ptr(raa) // OK 35 {} 36 #pragma omp target teams is_device_ptr(g) // OK 37 {} 38 #pragma omp target teams is_device_ptr(rg) // OK 39 {} 40 #pragma omp target teams is_device_ptr(da) // OK 41 {} 42 return; 43 } 44 }; 45 // CHECK: struct SA 46 // CHECK-NEXT: const int da[5] = {0}; 47 // CHECK-NEXT: ST g[10]; 48 // CHECK-NEXT: STarr &rg = this->g; 49 // CHECK-NEXT: int i; 50 // CHECK-NEXT: int &j = this->i; 51 // CHECK-NEXT: int *k = &this->j; 52 // CHECK-NEXT: int *&z = this->k; 53 // CHECK-NEXT: int aa[10]; 54 // CHECK-NEXT: arr &raa = this->aa; 55 // CHECK-NEXT: func( 56 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->k){{$}} 57 // CHECK-NEXT: { 58 // CHECK-NEXT: } 59 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->z) 60 // CHECK-NEXT: { 61 // CHECK-NEXT: } 62 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->aa) 63 // CHECK-NEXT: { 64 // CHECK-NEXT: } 65 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->raa) 66 // CHECK-NEXT: { 67 // CHECK-NEXT: } 68 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->g) 69 // CHECK-NEXT: { 70 // CHECK-NEXT: } 71 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->rg) 72 // CHECK-NEXT: { 73 // CHECK-NEXT: } 74 // CHECK-NEXT: #pragma omp target teams is_device_ptr(this->da) 75 76 struct SB { 77 unsigned A; 78 unsigned B; 79 float Arr[100]; 80 float *Ptr; fooSB81 float *foo() { 82 return &Arr[0]; 83 } 84 }; 85 86 struct SC { 87 unsigned A : 2; 88 unsigned B : 3; 89 unsigned C; 90 unsigned D; 91 float Arr[100]; 92 SB S; 93 SB ArrS[100]; 94 SB *PtrS; 95 SB *&RPtrS; 96 float *Ptr; 97 SCSC98 SC(SB *&_RPtrS) : RPtrS(_RPtrS) {} 99 }; 100 101 union SD { 102 unsigned A; 103 float B; 104 }; 105 106 struct S1; 107 extern S1 a; 108 class S2 { 109 mutable int a; 110 public: S2()111 S2():a(0) { } S2(S2 & s2)112 S2(S2 &s2):a(s2.a) { } 113 static float S2s; 114 static const float S2sc; 115 }; 116 const float S2::S2sc = 0; 117 const S2 b; 118 const S2 ba[5]; 119 class S3 { 120 int a; 121 public: S3()122 S3():a(0) { } S3(S3 & s3)123 S3(S3 &s3):a(s3.a) { } 124 }; 125 const S3 c; 126 const S3 ca[5]; 127 extern const int f; 128 class S4 { 129 int a; 130 S4(); 131 S4(const S4 &s4); 132 public: S4(int v)133 S4(int v):a(v) { } 134 }; 135 class S5 { 136 int a; S5()137 S5():a(0) {} S5(const S5 & s5)138 S5(const S5 &s5):a(s5.a) { } 139 public: S5(int v)140 S5(int v):a(v) { } 141 }; 142 143 S3 h; 144 #pragma omp threadprivate(h) 145 146 typedef struct { 147 int a; 148 } S6; 149 150 template <typename T> tmain(T argc)151T tmain(T argc) { 152 const T da[5] = { 0 }; 153 S6 h[10]; 154 auto &rh = h; 155 T i; 156 T &j = i; 157 T *k = &j; 158 T *&z = k; 159 T aa[10]; 160 auto &raa = aa; 161 #pragma omp target teams is_device_ptr(k) 162 {} 163 #pragma omp target teams is_device_ptr(z) 164 {} 165 #pragma omp target teams is_device_ptr(aa) 166 {} 167 #pragma omp target teams is_device_ptr(raa) 168 {} 169 #pragma omp target teams is_device_ptr(h) 170 {} 171 #pragma omp target teams is_device_ptr(rh) 172 {} 173 #pragma omp target teams is_device_ptr(da) 174 {} 175 return 0; 176 } 177 178 // CHECK: template<> int tmain<int>(int argc) { 179 // CHECK-NEXT: const int da[5] = {0}; 180 // CHECK-NEXT: S6 h[10]; 181 // CHECK-NEXT: auto &rh = h; 182 // CHECK-NEXT: int i; 183 // CHECK-NEXT: int &j = i; 184 // CHECK-NEXT: int *k = &j; 185 // CHECK-NEXT: int *&z = k; 186 // CHECK-NEXT: int aa[10]; 187 // CHECK-NEXT: auto &raa = aa; 188 // CHECK-NEXT: #pragma omp target teams is_device_ptr(k) 189 // CHECK-NEXT: { 190 // CHECK-NEXT: } 191 // CHECK-NEXT: #pragma omp target teams is_device_ptr(z) 192 // CHECK-NEXT: { 193 // CHECK-NEXT: } 194 // CHECK-NEXT: #pragma omp target teams is_device_ptr(aa) 195 // CHECK-NEXT: { 196 // CHECK-NEXT: } 197 // CHECK-NEXT: #pragma omp target teams is_device_ptr(raa) 198 // CHECK-NEXT: { 199 // CHECK-NEXT: } 200 // CHECK-NEXT: #pragma omp target teams is_device_ptr(h) 201 // CHECK-NEXT: { 202 // CHECK-NEXT: } 203 // CHECK-NEXT: #pragma omp target teams is_device_ptr(rh) 204 // CHECK-NEXT: { 205 // CHECK-NEXT: } 206 // CHECK-NEXT: #pragma omp target teams is_device_ptr(da) 207 208 // CHECK: template<> int *tmain<int *>(int *argc) { 209 // CHECK-NEXT: int *const da[5] = {0}; 210 // CHECK-NEXT: S6 h[10]; 211 // CHECK-NEXT: auto &rh = h; 212 // CHECK-NEXT: int *i; 213 // CHECK-NEXT: int *&j = i; 214 // CHECK-NEXT: int **k = &j; 215 // CHECK-NEXT: int **&z = k; 216 // CHECK-NEXT: int *aa[10]; 217 // CHECK-NEXT: auto &raa = aa; 218 // CHECK-NEXT: #pragma omp target teams is_device_ptr(k) 219 // CHECK-NEXT: { 220 // CHECK-NEXT: } 221 // CHECK-NEXT: #pragma omp target teams is_device_ptr(z) 222 // CHECK-NEXT: { 223 // CHECK-NEXT: } 224 // CHECK-NEXT: #pragma omp target teams is_device_ptr(aa) 225 // CHECK-NEXT: { 226 // CHECK-NEXT: } 227 // CHECK-NEXT: #pragma omp target teams is_device_ptr(raa) 228 // CHECK-NEXT: { 229 // CHECK-NEXT: } 230 // CHECK-NEXT: #pragma omp target teams is_device_ptr(h) 231 // CHECK-NEXT: { 232 // CHECK-NEXT: } 233 // CHECK-NEXT: #pragma omp target teams is_device_ptr(rh) 234 // CHECK-NEXT: { 235 // CHECK-NEXT: } 236 // CHECK-NEXT: #pragma omp target teams is_device_ptr(da) 237 238 // CHECK-LABEL: int main(int argc, char **argv) { main(int argc,char ** argv)239int main(int argc, char **argv) { 240 const int da[5] = { 0 }; 241 S6 h[10]; 242 auto &rh = h; 243 int i; 244 int &j = i; 245 int *k = &j; 246 int *&z = k; 247 int aa[10]; 248 auto &raa = aa; 249 // CHECK-NEXT: const int da[5] = {0}; 250 // CHECK-NEXT: S6 h[10]; 251 // CHECK-NEXT: auto &rh = h; 252 // CHECK-NEXT: int i; 253 // CHECK-NEXT: int &j = i; 254 // CHECK-NEXT: int *k = &j; 255 // CHECK-NEXT: int *&z = k; 256 // CHECK-NEXT: int aa[10]; 257 // CHECK-NEXT: auto &raa = aa; 258 #pragma omp target teams is_device_ptr(k) 259 // CHECK-NEXT: #pragma omp target teams is_device_ptr(k) 260 {} 261 // CHECK-NEXT: { 262 // CHECK-NEXT: } 263 #pragma omp target teams is_device_ptr(z) 264 // CHECK-NEXT: #pragma omp target teams is_device_ptr(z) 265 {} 266 // CHECK-NEXT: { 267 // CHECK-NEXT: } 268 #pragma omp target teams is_device_ptr(aa) 269 // CHECK-NEXT: #pragma omp target teams is_device_ptr(aa) 270 {} 271 // CHECK-NEXT: { 272 // CHECK-NEXT: } 273 #pragma omp target teams is_device_ptr(raa) 274 // CHECK-NEXT: #pragma omp target teams is_device_ptr(raa) 275 {} 276 // CHECK-NEXT: { 277 // CHECK-NEXT: } 278 #pragma omp target teams is_device_ptr(h) 279 // CHECK-NEXT: #pragma omp target teams is_device_ptr(h) 280 {} 281 // CHECK-NEXT: { 282 // CHECK-NEXT: } 283 #pragma omp target teams is_device_ptr(rh) 284 // CHECK-NEXT: #pragma omp target teams is_device_ptr(rh) 285 {} 286 // CHECK-NEXT: { 287 // CHECK-NEXT: } 288 #pragma omp target teams is_device_ptr(da) 289 // CHECK-NEXT: #pragma omp target teams is_device_ptr(da) 290 {} 291 // CHECK-NEXT: { 292 // CHECK-NEXT: } 293 return tmain<int>(argc) + *tmain<int *>(&argc); 294 } 295 #endif 296