clang/test/OpenMP/target_has_device_addr_ast_print.cpp

   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 -fsyntax-only -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 -fsyntax-only -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;
  39   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;
  93   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
 110   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:
 123   S2():a(0) { }
 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:
 134   S3():a(0) { }
 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:
 145   S4(int v):a(v) { }
 146 };
 147 class S5 {
 148   int a;
 149   S5():a(0) {}
 150   S5(const S5 &s5):a(s5.a) { }
 151 public:
 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>
 163 T 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) {
 251 int 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>
 320   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
 334 void use_template() {
 335   SomeKernel aKern;
 336   aKern.apply<32>();
 337 }
 338 #endif