clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
   7 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
   8 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-64
   9 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
  10 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  11 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
  12
  13 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  14 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  15 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  16 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  17 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  18 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  19 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
  20 #ifdef CK1
  21
  22 double *g;
  23
  24 // CK1: @g ={{.*}} global ptr
  25 // CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 19, i64 64]
  26 // CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 67]
  27 // CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 67]
  28 // CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 67]
  29 // CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 67]
  30 // CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 67]
  31 // CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 67]
  32 // CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 3]
  33 // CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
  34 // CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
  35 // CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
  36 // CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
  37
  38 // CK1-LABEL: @_Z3foo
  39 template<typename T>
  40 void foo(float *&lr, T *&tr) {
  41   float *l;
  42   T *t;
  43
  44   // CK1:     [[T:%.+]] = load ptr, ptr [[DECL:@g]],
  45   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
  46   // CK1:     store ptr [[T]], ptr [[BP]],
  47   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
  48   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
  49   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
  50   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
  51   // CK1:     [[TT:%.+]] = load ptr, ptr [[PVT]],
  52   // CK1:     getelementptr inbounds double, ptr [[TT]], i32 1
  53   #pragma omp target data map(g[:10]) use_device_ptr(g)
  54   {
  55     ++g;
  56   }
  57   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
  58   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
  59   // CK1:     getelementptr inbounds double, ptr [[TTT]], i32 1
  60   ++g;
  61
  62   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
  63   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
  64   // CK1:     store ptr [[T1]], ptr [[BP]],
  65   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
  66   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
  67   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
  68   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
  69   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
  70   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
  71   #pragma omp target data map(l[:10]) use_device_ptr(l)
  72   {
  73     ++l;
  74   }
  75   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE01]]
  76   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
  77   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
  78   ++l;
  79
  80   // CK1-NOT: call void @__tgt_target
  81   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
  82   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
  83   #pragma omp target data map(l[:10]) use_device_ptr(l) if(0)
  84   {
  85     ++l;
  86   }
  87   // CK1-NOT: call void @__tgt_target
  88   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
  89   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
  90   ++l;
  91
  92   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
  93   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
  94   // CK1:     store ptr [[T1]], ptr [[BP]],
  95   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
  96   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
  97   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
  98   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
  99   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 100   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
 101   #pragma omp target data map(l[:10]) use_device_ptr(l) if(1)
 102   {
 103     ++l;
 104   }
 105   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE03]]
 106   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 107   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
 108   ++l;
 109
 110   // CK1:     [[CMP:%.+]] = icmp ne ptr %{{.+}}, null
 111   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
 112
 113   // CK1:     [[BTHEN]]:
 114   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
 115   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
 116   // CK1:     store ptr [[T1]], ptr [[BP]],
 117   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE04]]
 118   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 119   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
 120   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
 121   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 122   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
 123   // CK1:     br label %[[BEND:.+]]
 124
 125   // CK1:     [[BELSE]]:
 126   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 127   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
 128   // CK1:     br label %[[BEND]]
 129   #pragma omp target data map(l[:10]) use_device_ptr(l) if(lr != 0)
 130   {
 131     ++l;
 132   }
 133   // CK1:     [[BEND]]:
 134   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
 135
 136   // CK1:     [[BTHEN]]:
 137   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE04]]
 138   // CK1:     br label %[[BEND:.+]]
 139
 140   // CK1:     [[BELSE]]:
 141   // CK1:     br label %[[BEND]]
 142
 143   // CK1:     [[BEND]]:
 144   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 145   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
 146   ++l;
 147
 148   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
 149   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
 150   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
 151   // CK1:     store ptr [[T1]], ptr [[BP]],
 152   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE05]]
 153   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 154   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
 155   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
 156   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
 157   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 158   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
 159   // CK1:     getelementptr inbounds float, ptr [[TT2]], i32 1
 160   #pragma omp target data map(lr[:10]) use_device_ptr(lr)
 161   {
 162     ++lr;
 163   }
 164   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE05]]
 165   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 166   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
 167   // CK1:     getelementptr inbounds float, ptr [[TTTT]], i32 1
 168   ++lr;
 169
 170   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
 171   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
 172   // CK1:     store ptr [[T1]], ptr [[BP]],
 173   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE06]]
 174   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 175   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
 176   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
 177   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 178   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
 179   #pragma omp target data map(t[:10]) use_device_ptr(t)
 180   {
 181     ++t;
 182   }
 183   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE06]]
 184   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 185   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
 186   ++t;
 187
 188   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
 189   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
 190   // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
 191   // CK1:     store ptr [[T1]], ptr [[BP]],
 192   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE07]]
 193   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 194   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
 195   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
 196   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
 197   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 198   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
 199   // CK1:     getelementptr inbounds i32, ptr [[TT2]], i32 1
 200   #pragma omp target data map(tr[:10]) use_device_ptr(tr)
 201   {
 202     ++tr;
 203   }
 204   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE07]]
 205   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 206   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
 207   // CK1:     getelementptr inbounds i32, ptr [[TTTT]], i32 1
 208   ++tr;
 209
 210   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
 211   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 0
 212   // CK1:     store ptr [[T1]], ptr [[BP]],
 213   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE08]]
 214   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 215   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
 216   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
 217   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 218   // CK1:     getelementptr inbounds float, ptr [[TT1]], i32 1
 219   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l)
 220   {
 221     ++l; ++t;
 222   }
 223   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE08]]
 224   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 225   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
 226   ++l; ++t;
 227
 228
 229   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE09]]
 230   // CK1:     [[_VAL:%.+]] = load ptr, ptr {{%.+}},
 231   // CK1:     store ptr [[_VAL]], ptr [[_PVT:%.+]],
 232   // CK1:     [[VAL:%.+]] = load ptr, ptr {{%.+}},
 233   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
 234   // CK1:     [[_TT1:%.+]] = load ptr, ptr [[_PVT]],
 235   // CK1:     getelementptr inbounds float, ptr [[_TT1]], i32 1
 236   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 237   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
 238   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l) use_device_ptr(t)
 239   {
 240     ++l; ++t;
 241   }
 242   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE09]]
 243   // CK1:     [[_TTT:%.+]] = load ptr, ptr {{%.+}},
 244   // CK1:     getelementptr inbounds float, ptr [[_TTT]], i32 1
 245   // CK1:     [[TTT:%.+]] = load ptr, ptr {{%.+}},
 246   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
 247   ++l; ++t;
 248
 249   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE10]]
 250   // CK1:     [[_VAL:%.+]] = load ptr, ptr {{%.+}},
 251   // CK1:     store ptr [[_VAL]], ptr [[_PVT:%.+]],
 252   // CK1:     [[VAL:%.+]] = load ptr, ptr {{%.+}},
 253   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
 254   // CK1:     [[_TT1:%.+]] = load ptr, ptr [[_PVT]],
 255   // CK1:     getelementptr inbounds float, ptr [[_TT1]], i32 1
 256   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 257   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
 258   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l,t)
 259   {
 260     ++l; ++t;
 261   }
 262   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE10]]
 263   // CK1:     [[_TTT:%.+]] = load ptr, ptr {{%.+}},
 264   // CK1:     getelementptr inbounds float, ptr [[_TTT]], i32 1
 265   // CK1:     [[TTT:%.+]] = load ptr, ptr {{%.+}},
 266   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
 267   ++l; ++t;
 268
 269   // CK1:     [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],
 270   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
 271   // CK1:     store ptr [[T1]], ptr [[BP]],
 272   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE11]]
 273   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 274   // CK1-NOT: store ptr [[VAL]], ptr [[DECL]],
 275   // CK1:     store ptr [[VAL]], ptr [[PVT:%.+]],
 276   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 277   // CK1:     getelementptr inbounds i32, ptr [[TT1]], i32 1
 278   #pragma omp target data map(l[:10]) use_device_ptr(t)
 279   {
 280     ++l; ++t;
 281   }
 282   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE11]]
 283   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 284   // CK1:     getelementptr inbounds i32, ptr [[TTT]], i32 1
 285   ++l; ++t;
 286
 287   // CK1:     [[T2:%.+]] = load ptr, ptr [[DECL:%.+]],
 288   // CK1:     [[T1:%.+]] = load ptr, ptr [[T2]],
 289   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
 290   // CK1:     store ptr [[T1]], ptr [[BP]],
 291   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE12]]
 292   // CK1:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 293   // CK1:     store ptr [[VAL]], ptr [[PVTV:%.+]],
 294   // CK1-NOT: store ptr [[PVTV]], ptr [[DECL]],
 295   // CK1:     store ptr [[PVTV]], ptr [[PVT:%.+]],
 296   // CK1:     [[TT1:%.+]] = load ptr, ptr [[PVT]],
 297   // CK1:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
 298   // CK1:     getelementptr inbounds i32, ptr [[TT2]], i32 1
 299   #pragma omp target data map(l[:10]) use_device_ptr(tr)
 300   {
 301     ++l; ++tr;
 302   }
 303   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE12]]
 304   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 305   // CK1:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
 306   // CK1:     getelementptr inbounds i32, ptr [[TTTT]], i32 1
 307   ++l; ++tr;
 308
 309 }
 310
 311 void bar(float *&a, int *&b) {
 312   foo<int>(a,b);
 313 }
 314
 315 #endif
 316 ///==========================================================================///
 317 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
 318 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 319 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
 320 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 321 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 322 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 323
 324 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 325 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 326 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 327 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 328 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 329 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 330 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
 331 #ifdef CK2
 332
 333 // CK2: [[ST:%.+]] = type { ptr, ptr }
 334 // CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
 335 // CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
 336 // CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 3, i64 0, i64 562949953421392]
 337 // CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 0, i64 281474976710739, i64 281474976710736]
 338
 339 template <typename T>
 340 struct ST {
 341   T *a;
 342   double *&b;
 343   ST(double *&b) : a(0), b(b) {}
 344
 345   // CK2-LABEL: @{{.*}}foo{{.*}}
 346   void foo(double *&arg) {
 347     int *la = 0;
 348
 349     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
 350     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
 351     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
 352     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 353     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
 354     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
 355     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
 356     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
 357     // CK2:     getelementptr inbounds double, ptr [[TT2]], i32 1
 358     #pragma omp target data map(a[:10]) use_device_ptr(a)
 359     {
 360       a++;
 361     }
 362     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
 363     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 0
 364     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 365     // CK2:     getelementptr inbounds double, ptr [[TTT]], i32 1
 366     a++;
 367
 368     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1
 369     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
 370     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
 371     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 372     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
 373     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
 374     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
 375     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
 376     // CK2:     getelementptr inbounds double, ptr [[TT2]], i32 1
 377     #pragma omp target data map(b[:10]) use_device_ptr(b)
 378     {
 379       b++;
 380     }
 381     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE01]]
 382     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %{{.+}}, i32 0, i32 1
 383     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 384     // CK2:     [[TTTT:%.+]] = load ptr, ptr [[TTT]],
 385     // CK2:     getelementptr inbounds double, ptr [[TTTT]], i32 1
 386     b++;
 387
 388     // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
 389     // CK2:     store ptr [[RVAL:%.+]], ptr [[BP]],
 390     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]
 391     // CK2:     [[VAL:%.+]] = load ptr, ptr [[BP]],
 392     // CK2:     store ptr [[VAL]], ptr [[PVT:%.+]],
 393     // CK2:     store ptr [[PVT]], ptr [[PVT2:%.+]],
 394     // CK2:     [[TT1:%.+]] = load ptr, ptr [[PVT2]],
 395     // CK2:     [[TT2:%.+]] = load ptr, ptr [[TT1]],
 396     // CK2:     getelementptr inbounds double, ptr [[TT2]], i32 1
 397     #pragma omp target data map(la[:10]) use_device_ptr(a)
 398     {
 399       a++;
 400       la++;
 401     }
 402     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE02]]
 403     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 0
 404     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 405     // CK2:     getelementptr inbounds double, ptr [[TTT]], i32 1
 406     a++;
 407     la++;
 408
 409     // CK2:     [[BP1:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 1
 410     // CK2:     store ptr [[RVAL1:%.+]], ptr [[BP1]],
 411     // CK2:     [[BP2:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
 412     // CK2:     store ptr [[RVAL2:%.+]], ptr [[BP2]],
 413     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
 414     // CK2:     [[VAL1:%.+]] = load ptr, ptr [[BP1]],
 415     // CK2:     store ptr [[VAL1]], ptr [[PVT1:%.+]],
 416     // CK2:     [[VAL2:%.+]] = load ptr, ptr [[BP2]],
 417     // CK2:     store ptr [[VAL2]], ptr [[PVT2:%.+]],
 418     // CK2:     store ptr [[PVT2]], ptr [[_PVT2:%.+]],
 419     // CK2:     store ptr [[PVT1]], ptr [[_PVT1:%.+]],
 420     // CK2:     [[TT2:%.+]] = load ptr, ptr [[_PVT2]],
 421     // CK2:     [[_TT2:%.+]] = load ptr, ptr [[TT2]],
 422     // CK2:     getelementptr inbounds double, ptr [[_TT2]], i32 1
 423     // CK2:     [[TT1:%.+]] = load ptr, ptr [[_PVT1]],
 424     // CK2:     [[_TT1:%.+]] = load ptr, ptr [[TT1]],
 425     // CK2:     getelementptr inbounds double, ptr [[_TT1]], i32 1
 426     #pragma omp target data map(b[:10]) use_device_ptr(a, b)
 427     {
 428       a++;
 429       b++;
 430     }
 431     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE03]]
 432     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 0
 433     // CK2:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
 434     // CK2:     getelementptr inbounds double, ptr [[TTT]], i32 1
 435     // CK2:     [[_DECL:%.+]] = getelementptr inbounds [[ST]], ptr %this1, i32 0, i32 1
 436     // CK2:     [[_TTT:%.+]] = load ptr, ptr [[_DECL]],
 437     // CK2:     [[_TTTT:%.+]] = load ptr, ptr [[_TTT]],
 438     // CK2:     getelementptr inbounds double, ptr [[_TTTT]], i32 1
 439     a++;
 440     b++;
 441   }
 442 };
 443
 444 void bar(double *arg){
 445   ST<double> A(arg);
 446   A.foo(arg);
 447   ++arg;
 448 }
 449 #endif
 450 #endif