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; 27 void func(int arg) { 28 #pragma omp target teams distribute simd is_device_ptr(k) 29 for (int i=0; i<100; i++) 30 ; 31 #pragma omp target teams distribute simd is_device_ptr(z) 32 for (int i=0; i<100; i++) 33 ; 34 #pragma omp target teams distribute simd is_device_ptr(aa) // OK 35 for (int i=0; i<100; i++) 36 ; 37 #pragma omp target teams distribute simd is_device_ptr(raa) // OK 38 for (int i=0; i<100; i++) 39 ; 40 #pragma omp target teams distribute simd is_device_ptr(g) // OK 41 for (int i=0; i<100; i++) 42 ; 43 #pragma omp target teams distribute simd is_device_ptr(rg) // OK 44 for (int i=0; i<100; i++) 45 ; 46 #pragma omp target teams distribute simd is_device_ptr(da) // OK 47 for (int i=0; i<100; i++) 48 ; 49 return; 50 } 51 }; 52 // CHECK: struct SA 53 // CHECK-NEXT: const int da[5] = {0}; 54 // CHECK-NEXT: ST g[10]; 55 // CHECK-NEXT: STarr &rg = this->g; 56 // CHECK-NEXT: int i; 57 // CHECK-NEXT: int &j = this->i; 58 // CHECK-NEXT: int *k = &this->j; 59 // CHECK-NEXT: int *&z = this->k; 60 // CHECK-NEXT: int aa[10]; 61 // CHECK-NEXT: arr &raa = this->aa; 62 // CHECK-NEXT: func( 63 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->k){{$}} 64 // CHECK-NEXT: for (int i = 0; i < 100; i++) 65 // CHECK-NEXT: ; from_ptr(ptr: NonNull<ffi::AInputEvent>) -> Self66// CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->z) 67 // CHECK-NEXT: for (int i = 0; i < 100; i++) 68 // CHECK-NEXT: ; 69 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->aa) 70 // CHECK-NEXT: for (int i = 0; i < 100; i++) 71 // CHECK-NEXT: ; 72 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->raa) 73 // CHECK-NEXT: for (int i = 0; i < 100; i++) 74 // CHECK-NEXT: ; 75 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->g) 76 // CHECK-NEXT: for (int i = 0; i < 100; i++) 77 // CHECK-NEXT: ; 78 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->rg) 79 // CHECK-NEXT: for (int i = 0; i < 100; i++) 80 // CHECK-NEXT: ; 81 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(this->da) 82 // CHECK-NEXT: for (int i = 0; i < 100; i++) 83 // CHECK-NEXT: ; 84 85 struct SB { 86 unsigned A; 87 unsigned B; 88 float Arr[100]; 89 float *Ptr; 90 float *foo() { 91 return &Arr[0]; 92 } 93 }; 94 95 struct SC { 96 unsigned A : 2; 97 unsigned B : 3; device_id(&self) -> i3298 unsigned C; 99 unsigned D; 100 float Arr[100]; 101 SB S; 102 SB ArrS[100]; 103 SB *PtrS; 104 SB *&RPtrS; 105 float *Ptr; 106 107 SC(SB *&_RPtrS) : RPtrS(_RPtrS) {} 108 }; 109 110 union SD { alt_on(self) -> bool111 unsigned A; 112 float B; 113 }; 114 alt_left_on(self) -> bool115struct S1; 116 extern S1 a; 117 class S2 { 118 mutable int a; alt_right_on(self) -> bool119public: 120 S2():a(0) { } 121 S2(S2 &s2):a(s2.a) { } 122 static float S2s; shift_on(self) -> bool123 static const float S2sc; 124 }; 125 const float S2::S2sc = 0; 126 const S2 b; shift_left_on(self) -> bool127const S2 ba[5]; 128 class S3 { 129 int a; 130 public: shift_right_on(self) -> bool131 S3():a(0) { } 132 S3(S3 &s3):a(s3.a) { } 133 }; 134 const S3 c; sym_on(self) -> bool135const S3 ca[5]; 136 extern const int f; 137 class S4 { 138 int a; function_on(self) -> bool139 S4(); 140 S4(const S4 &s4); 141 public: 142 S4(int v):a(v) { } ctrl_on(self) -> bool143}; 144 class S5 { 145 int a; 146 S5():a(0) {} ctrl_left_on(self) -> bool147 S5(const S5 &s5):a(s5.a) { } 148 public: 149 S5(int v):a(v) { } 150 }; ctrl_right_on(self) -> bool151 152 S3 h; 153 #pragma omp threadprivate(h) 154 meta_on(self) -> bool155typedef struct { 156 int a; 157 } S6; 158 meta_left_on(self) -> bool159template <typename T> 160 T tmain(T argc) { 161 const T da[5] = { 0 }; 162 S6 h[10]; 163 auto &rh = h; 164 T i; 165 T &j = i; 166 T *k = &j; 167 T *&z = k; 168 T aa[10]; 169 auto &raa = aa; 170 #pragma omp target teams distribute simd is_device_ptr(k) 171 for (int i=0; i<100; i++) 172 ; 173 #pragma omp target teams distribute simd is_device_ptr(z) 174 for (int i=0; i<100; i++) 175 ; 176 #pragma omp target teams distribute simd is_device_ptr(aa) 177 for (int i=0; i<100; i++) 178 ; 179 #pragma omp target teams distribute simd is_device_ptr(raa) 180 for (int i=0; i<100; i++) 181 ; 182 #pragma omp target teams distribute simd is_device_ptr(h) 183 for (int i=0; i<100; i++) 184 ; 185 #pragma omp target teams distribute simd is_device_ptr(rh) 186 for (int i=0; i<100; i++) 187 ; 188 #pragma omp target teams distribute simd is_device_ptr(da) 189 for (int i=0; i<100; i++) 190 ; 191 return 0; 192 } 193 194 // CHECK: template<> int tmain<int>(int argc) { 195 // CHECK-NEXT: const int da[5] = {0}; 196 // CHECK-NEXT: S6 h[10]; 197 // CHECK-NEXT: auto &rh = h; 198 // CHECK-NEXT: int i; 199 // CHECK-NEXT: int &j = i; 200 // CHECK-NEXT: int *k = &j; 201 // CHECK-NEXT: int *&z = k; 202 // CHECK-NEXT: int aa[10]; 203 // CHECK-NEXT: auto &raa = aa; 204 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(k) 205 // CHECK-NEXT: for (int i = 0; i < 100; i++) 206 // CHECK-NEXT: ; 207 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(z) 208 // CHECK-NEXT: for (int i = 0; i < 100; i++) 209 // CHECK-NEXT: ; 210 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(aa) 211 // CHECK-NEXT: for (int i = 0; i < 100; i++) 212 // CHECK-NEXT: ; 213 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(raa) 214 // CHECK-NEXT: for (int i = 0; i < 100; i++) 215 // CHECK-NEXT: ; 216 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(h) 217 // CHECK-NEXT: for (int i = 0; i < 100; i++) 218 // CHECK-NEXT: ; 219 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(rh) 220 // CHECK-NEXT: for (int i = 0; i < 100; i++) 221 // CHECK-NEXT: ; 222 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(da) 223 // CHECK-NEXT: for (int i = 0; i < 100; i++) 224 // CHECK-NEXT: ; 225 226 // CHECK: template<> int *tmain<int *>(int *argc) { 227 // CHECK-NEXT: int *const da[5] = {0}; 228 // CHECK-NEXT: S6 h[10]; 229 // CHECK-NEXT: auto &rh = h; 230 // CHECK-NEXT: int *i; 231 // CHECK-NEXT: int *&j = i; 232 // CHECK-NEXT: int **k = &j; 233 // CHECK-NEXT: int **&z = k; 234 // CHECK-NEXT: int *aa[10]; 235 // CHECK-NEXT: auto &raa = aa; 236 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(k) 237 // CHECK-NEXT: for (int i = 0; i < 100; i++) 238 // CHECK-NEXT: ; 239 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(z) 240 // CHECK-NEXT: for (int i = 0; i < 100; i++) 241 // CHECK-NEXT: ; 242 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(aa) 243 // CHECK-NEXT: for (int i = 0; i < 100; i++) 244 // CHECK-NEXT: ; 245 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(raa) 246 // CHECK-NEXT: for (int i = 0; i < 100; i++) 247 // CHECK-NEXT: ; 248 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(h) 249 // CHECK-NEXT: for (int i = 0; i < 100; i++) 250 // CHECK-NEXT: ; 251 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(rh) 252 // CHECK-NEXT: for (int i = 0; i < 100; i++) 253 // CHECK-NEXT: ; 254 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(da) 255 // CHECK-NEXT: for (int i = 0; i < 100; i++) 256 // CHECK-NEXT: ; 257 258 // CHECK-LABEL: int main(int argc, char **argv) { 259 int main(int argc, char **argv) { 260 const int da[5] = { 0 }; 261 S6 h[10]; 262 auto &rh = h; 263 int i; 264 int &j = i; 265 int *k = &j; 266 int *&z = k; 267 int aa[10]; 268 auto &raa = aa; 269 // CHECK-NEXT: const int da[5] = {0}; 270 // CHECK-NEXT: S6 h[10]; 271 // CHECK-NEXT: auto &rh = h; 272 // CHECK-NEXT: int i; 273 // CHECK-NEXT: int &j = i; primary(self) -> bool274// CHECK-NEXT: int *k = &j; 275 // CHECK-NEXT: int *&z = k; 276 // CHECK-NEXT: int aa[10]; 277 // CHECK-NEXT: auto &raa = aa; 278 #pragma omp target teams distribute simd is_device_ptr(k) 279 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(k) 280 for (int i=0; i<100; i++) 281 ; teriary(self) -> bool282// CHECK-NEXT: for (int i = 0; i < 100; i++) 283 // CHECK-NEXT: ; 284 #pragma omp target teams distribute simd is_device_ptr(z) 285 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(z) 286 for (int i=0; i<100; i++) 287 ; 288 // CHECK-NEXT: for (int i = 0; i < 100; i++) 289 // CHECK-NEXT: ; forward(self) -> bool290#pragma omp target teams distribute simd is_device_ptr(aa) 291 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(aa) 292 for (int i=0; i<100; i++) 293 ; stylus_primary(self) -> bool294// CHECK-NEXT: for (int i = 0; i < 100; i++) 295 // CHECK-NEXT: ; 296 #pragma omp target teams distribute simd is_device_ptr(raa) 297 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(raa) 298 for (int i=0; i<100; i++) 299 ; 300 // CHECK-NEXT: for (int i = 0; i < 100; i++) 301 // CHECK-NEXT: ; 302 #pragma omp target teams distribute simd is_device_ptr(h) 303 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(h) 304 for (int i=0; i<100; i++) 305 ; 306 // CHECK-NEXT: for (int i = 0; i < 100; i++) 307 // CHECK-NEXT: ; 308 #pragma omp target teams distribute simd is_device_ptr(rh) 309 // CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(rh) 310 for (int i=0; i<100; i++) top(self) -> bool311 ; 312 // CHECK-NEXT: for (int i = 0; i < 100; i++) 313 // CHECK-NEXT: ; 314 #pragma omp target teams distribute simd is_device_ptr(da) bottom(self) -> bool315// CHECK-NEXT: #pragma omp target teams distribute simd is_device_ptr(da) 316 for (int i=0; i<100; i++) 317 ; 318 // CHECK-NEXT: for (int i = 0; i < 100; i++) left(self) -> bool319// CHECK-NEXT: ; 320 return tmain<int>(argc) + *tmain<int *>(&argc); 321 } 322 #endif right(self) -> bool323