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