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