clang/test/OpenMP/target_data_codegen.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -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 -Wno-vla %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-64
   9 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -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 -Wno-vla %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
  12
  13 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -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 -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  16 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -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 -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  19 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
  20 #ifdef CK1
  21
  22 // CK1: [[ST:%.+]] = type { i32, ptr }
  23 template <typename T>
  24 struct ST {
  25   T a;
  26   double *b;
  27 };
  28
  29 ST<int> gb;
  30 double gc[100];
  31
  32 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
  33 // CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  34
  35 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
  36 // CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
  37
  38 // CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 5]
  39
  40 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24]
  41 // CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
  42
  43 // CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1025]
  44
  45 // CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1029]
  46
  47 // CK1-LABEL: _Z3fooi
  48 void foo(int arg) {
  49   int la;
  50   float lb[arg];
  51
  52   // Region 00
  53   // CK1-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE00]], ptr [[MTYPE00]], ptr null, ptr null)
  54   // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
  55   // CK1-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
  56   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  57   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  58
  59   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
  60   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
  61   // CK1-DAG: store ptr @gc, ptr [[BP0]]
  62   // CK1-DAG: store ptr @gc, ptr [[P0]]
  63
  64   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
  65
  66   // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE00]], ptr [[MTYPE00]], ptr null, ptr null)
  67   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
  68 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
  69   #pragma omp target data if(1+3-5) device(arg) map(from: gc)
  70   {++arg;}
  71
  72   // Region 01
  73   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
  74   #pragma omp target data map(la) if(1+3-4)
  75   {++arg;}
  76
  77   // Region 02
  78   // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
  79   // CK1: [[IFTHEN]]
  80   // CK1-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 4, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE02]], ptr [[MTYPE02]], ptr null, ptr null)
  81   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  82   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  83
  84   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
  85   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
  86   // CK1-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
  87   // CK1-DAG: store ptr [[VAR0]], ptr [[P0]]
  88   // CK1: br label %[[IFEND:[^,]+]]
  89
  90   // CK1: [[IFELSE]]
  91   // CK1: br label %[[IFEND]]
  92   // CK1: [[IFEND]]
  93   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
  94   // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
  95
  96   // CK1: [[IFTHEN]]
  97   // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 4, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE02]], ptr [[MTYPE02]], ptr null, ptr null)
  98   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
  99   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 100   // CK1: br label %[[IFEND:[^,]+]]
 101   // CK1: [[IFELSE]]
 102   // CK1: br label %[[IFEND]]
 103   // CK1: [[IFEND]]
 104   #pragma omp target data map(to: arg) if(arg) device(4)
 105   {++arg;}
 106
 107   // Region 03
 108   // CK1-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE03]], ptr null, ptr null)
 109   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 110   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 111   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 112
 113   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 114   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 115   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 116   // CK1-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 117   // CK1-DAG: store ptr [[VAR0]], ptr [[P0]]
 118   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], ptr [[S0]]
 119   // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
 120   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
 121   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
 122   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 123
 124   // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE03]], ptr null, ptr null)
 125   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 126   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 127   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
 128   #pragma omp target data map(always, to: lb)
 129   {++arg;}
 130
 131   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 132   {++arg;}
 133
 134   // Region 04
 135   // CK1-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE04]], ptr null, ptr null)
 136   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 137   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 138   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PSZ:%[^,]+]]
 139
 140   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 141   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 142   // CK1-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PSZ]], i{{.+}} 0, i{{.+}} 0
 143   // CK1-DAG: store ptr @gb, ptr [[BP0]]
 144   // CK1-DAG: store ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1), ptr [[P0]]
 145   // CK1-DAG: [[DIV:%.+]] = sdiv exact i64 sub (i64 ptrtoint (ptr getelementptr (ptr, ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1) to i64)), ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 146   // CK1-DAG: store i64 [[DIV]], ptr [[PS0]],
 147
 148   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 149   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 150   // CK1-DAG: store ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1), ptr [[BP1]]
 151   // CK1-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 152   // CK1-DAG: [[SEC1]] = getelementptr inbounds {{.+}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
 153   // CK1-DAG: [[SEC11]] = load ptr, ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1),
 154
 155   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 156
 157   // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE04]], ptr null, ptr null)
 158   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 159   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 160   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PSZ]]
 161   #pragma omp target data map(to: gb.b[:3])
 162   {++arg;}
 163
 164   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 165   {++arg;}
 166
 167   // Region 05
 168   // CK1-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE05]], ptr null, ptr null)
 169   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 170   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 171   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 172
 173   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 174   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 175   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 176   // CK1-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 177   // CK1-DAG: store ptr [[VAR0]], ptr [[P0]]
 178   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], ptr [[S0]]
 179   // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
 180   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
 181   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
 182   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 183
 184   // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE05]], ptr null, ptr null)
 185   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 186   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 187   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
 188   #pragma omp target data map(close, to: lb)
 189   {++arg;}
 190
 191   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 192   {++arg;}
 193
 194   // Region 06
 195   // CK1-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE06]], ptr null, ptr null)
 196   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 197   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 198   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 199
 200   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 201   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 202   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 203   // CK1-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 204   // CK1-DAG: store ptr [[VAR0]], ptr [[P0]]
 205   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], ptr [[S0]]
 206   // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
 207   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
 208   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
 209   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 210
 211   // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE06]], ptr null, ptr null)
 212   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 213   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 214   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
 215   #pragma omp target data map(always close, to: lb)
 216   {++arg;}
 217
 218 }
 219 #endif
 220 ///==========================================================================///
 221 // RUN: %clang_cc1 -DCK1A -verify -Wno-vla  -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
 222 // RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 223 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-64
 224 // RUN: %clang_cc1 -DCK1A -verify -Wno-vla  -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-32
 225 // RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 226 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-32
 227
 228 // RUN: %clang_cc1 -DCK1A -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 229 // RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 230 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 231 // RUN: %clang_cc1 -DCK1A -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 232 // RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 233 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 234 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 235 #ifdef CK1A
 236
 237 // CK1A: [[ST:%.+]] = type { i32, ptr }
 238 template <typename T>
 239 struct ST {
 240   T a;
 241   double *b;
 242 };
 243
 244 ST<int> gb;
 245 double gc[100];
 246
 247 // PRESENT=0x1000 | TO=0x1 = 0x1001
 248 // CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
 249
 250 // TO=0x1 = 0x1
 251 // CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1]]]
 252
 253 // PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405
 254 // CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
 255
 256 // CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x405
 257 // CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x405]]]
 258
 259 // CK1A-LABEL: _Z3fooi
 260 void foo(int arg) {
 261   int la;
 262   float lb[arg];
 263
 264   // Region 00
 265   // CK1A-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00Begin]]{{.+}})
 266   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 267   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 268   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 269
 270   // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 271   // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 272   // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 273   // CK1A-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 274   // CK1A-DAG: store ptr [[VAR0]], ptr [[P0]]
 275   // CK1A-DAG: store i[[sz:32|64]] [[CSVAL0:%[^,]+]], ptr [[S0]]
 276   // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
 277   // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
 278   // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
 279   // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 280
 281   // CK1A-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00End]]{{.+}})
 282   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 283   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 284   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
 285   #pragma omp target data map(present, to: lb)
 286   {++arg;}
 287
 288   // Region 01
 289   // CK1A-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE01Begin]]{{.+}})
 290   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 291   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 292   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 293
 294   // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 295   // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 296   // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 297   // CK1A-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 298   // CK1A-DAG: store ptr [[VAR0]], ptr [[P0]]
 299   // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], ptr [[S0]]
 300   // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
 301   // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
 302   // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
 303   // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 304
 305   // CK1A-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE01End]]{{.+}})
 306   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 307   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 308   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
 309   #pragma omp target data map(always close present, to: lb)
 310   {++arg;}
 311
 312 }
 313 #endif
 314 ///==========================================================================///
 315 // RUN: %clang_cc1 -DCK2 -verify -Wno-vla  -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
 316 // 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
 317 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
 318 // RUN: %clang_cc1 -DCK2 -verify -Wno-vla  -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
 319 // 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
 320 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 321
 322 // RUN: %clang_cc1 -DCK2 -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 323 // 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
 324 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 325 // RUN: %clang_cc1 -DCK2 -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 326 // 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
 327 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 328 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
 329 #ifdef CK2
 330
 331 // CK2: [[ST:%.+]] = type { i32, ptr }
 332 template <typename T>
 333 struct ST {
 334   T a;
 335   double *b;
 336
 337   T foo(T arg) {
 338     // Region 00
 339     #pragma omp target data map(always, to: b[1:3]) if(a>123) device(arg)
 340     {arg++;}
 341     return arg;
 342   }
 343 };
 344
 345 // CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24]
 346 // CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710677]
 347
 348 // CK2-LABEL: _Z3bari
 349 int bar(int arg){
 350   ST<int> A;
 351   return A.foo(arg);
 352 }
 353
 354 // Region 00
 355 // CK2-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64
 356 // CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
 357 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 358 // CK2: [[IFTHEN]]
 359 // CK2-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
 360 // CK2-DAG: [[GEPBP]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]]
 361 // CK2-DAG: [[GEPP]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]]
 362 // CK2-DAG: [[GEPS]] = getelementptr inbounds [2 x i64], ptr [[PS:%[^,]+]]
 363
 364 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 365 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 366 // CK2-DAG: [[PS0:%.+]] = getelementptr inbounds [2 x i64], ptr [[PS]], i{{.+}} 0, i{{.+}} 0
 367 // CK2-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 368 // CK2-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 369 // CK2-DAG: store i64 {{%.+}}, ptr [[PS0]],
 370 // CK2-DAG: [[SEC0]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1
 371
 372 // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 373 // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 374 // CK2-DAG: store ptr [[SEC0]], ptr [[BP1]]
 375 // CK2-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 376 // CK2-DAG: [[SEC1]] = getelementptr inbounds {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 1
 377 // CK2-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 378 // CK2-DAG: [[SEC111]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1
 379
 380 // CK2: br label %[[IFEND:[^,]+]]
 381
 382 // CK2: [[IFELSE]]
 383 // CK2: br label %[[IFEND]]
 384 // CK2: [[IFEND]]
 385 // CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 386 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 387
 388 // CK2: [[IFTHEN]]
 389 // CK2-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
 390 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 391 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 392 // CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
 393 // CK2: br label %[[IFEND:[^,]+]]
 394 // CK2: [[IFELSE]]
 395 // CK2: br label %[[IFEND]]
 396 // CK2: [[IFEND]]
 397 #endif
 398 ///==========================================================================///
 399 // RUN: %clang_cc1 -DCK3 -verify -Wno-vla  -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
 400 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 401 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-64
 402 // RUN: %clang_cc1 -DCK3 -verify -Wno-vla  -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
 403 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 404 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
 405
 406 // RUN: %clang_cc1 -DCK3 -verify -Wno-vla  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 407 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 408 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 409 // RUN: %clang_cc1 -DCK3 -verify -Wno-vla  -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 410 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 411 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 412 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 413 #ifdef CK3
 414
 415 // CK3-LABEL: no_target_devices
 416 void no_target_devices(int arg) {
 417   // CK3-NOT: tgt_target_data_begin
 418   // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 419   // CK3-NOT: tgt_target_data_end
 420   // CK3: ret
 421   #pragma omp target data map(to: arg) if(arg) device(4)
 422   {++arg;}
 423 }
 424 #endif
 425 ///==========================================================================///
 426 // RUN: %clang_cc1 -DCK4 -verify -Wno-vla  -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
 427 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 428 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK4 --check-prefix CK4-64
 429 // RUN: %clang_cc1 -DCK4 -verify -Wno-vla  -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK4 --check-prefix CK4-32
 430 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 431 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK4 --check-prefix CK4-32
 432
 433 // RUN: %clang_cc1 -DCK4 -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 434 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 435 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 436 // RUN: %clang_cc1 -DCK4 -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 437 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 438 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
 439 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
 440 #ifdef CK4
 441
 442 // CK4: [[STT:%.+]] = type { i32, ptr }
 443 template <typename T>
 444 struct STT {
 445   T a;
 446   double *b;
 447
 448   T foo(T arg) {
 449     // Region 00
 450     #pragma omp target data map(always, close to: b[1:3]) if(a>123) device(arg)
 451     {arg++;}
 452     return arg;
 453   }
 454 };
 455
 456 // CK4: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24]
 457 // CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711701]
 458
 459 // CK4-LABEL: _Z3bari
 460 int bar(int arg){
 461   STT<int> A;
 462   return A.foo(arg);
 463 }
 464
 465 // Region 00
 466 // CK4-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64
 467 // CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
 468 // CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 469 // CK4: [[IFTHEN]]
 470 // CK4-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
 471 // CK4-DAG: [[GEPBP]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]]
 472 // CK4-DAG: [[GEPP]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]]
 473 // CK4-DAG: [[GEPS]] = getelementptr inbounds [2 x i64], ptr [[PS:%[^,]+]]
 474
 475 // CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 476 // CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 477 // CK4-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i{{.+}} 0, i{{.+}} 0
 478 // CK4-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 479 // CK4-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 480 // CK4-DAG: store i64 {{%.+}}, ptr [[PS0]],
 481 // CK4-DAG: [[SEC0]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1
 482
 483 // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 484 // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 485 // CK4-DAG: store ptr [[SEC0]], ptr [[BP1]]
 486 // CK4-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 487 // CK4-DAG: [[SEC1]] = getelementptr inbounds {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 1
 488 // CK4-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 489 // CK4-DAG: [[SEC111]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1
 490
 491 // CK4: br label %[[IFEND:[^,]+]]
 492
 493 // CK4: [[IFELSE]]
 494 // CK4: br label %[[IFEND]]
 495 // CK4: [[IFEND]]
 496 // CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 497 // CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 498
 499 // CK4: [[IFTHEN]]
 500 // CK4-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
 501 // CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 502 // CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 503 // CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
 504 // CK4: br label %[[IFEND:[^,]+]]
 505 // CK4: [[IFELSE]]
 506 // CK4: br label %[[IFEND]]
 507 // CK4: [[IFEND]]
 508 #endif
 509 ///==========================================================================///
 510 // RUN: %clang_cc1 -DCK5 -verify -Wno-vla  -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
 511 // RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 512 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-64
 513 // RUN: %clang_cc1 -DCK5 -verify -Wno-vla  -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
 514 // RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 515 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
 516
 517 // RUN: %clang_cc1 -DCK5 -verify -Wno-vla  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 518 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 519 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 520 // RUN: %clang_cc1 -DCK5 -verify -Wno-vla  -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 521 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 522 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 523 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK5
 524 #ifdef CK5
 525 struct S1 {
 526   int i;
 527 };
 528 struct S2 {
 529   S1 s;
 530   struct S2 *ps;
 531 };
 532
 533 void test_close_modifier(int arg) {
 534   S2 *ps;
 535 // CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043]
 536 #pragma omp target data map(close, tofrom \
 537                             : arg, ps->ps->ps->ps->s)
 538   {
 539     ++(arg);
 540   }
 541 }
 542 #endif
 543 ///==========================================================================///
 544 // RUN: %clang_cc1 -DCK6 -verify -Wno-vla  -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64
 545 // RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 546 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-64
 547 // RUN: %clang_cc1 -DCK6 -verify -Wno-vla  -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
 548 // RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 549 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
 550
 551 // RUN: %clang_cc1 -DCK6 -verify -Wno-vla  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 552 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 553 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 554 // RUN: %clang_cc1 -DCK6 -verify -Wno-vla  -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 555 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 556 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 557 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 558 #ifdef CK6
 559 void test_close_modifier(int arg) {
 560 // CK6: private unnamed_addr constant [1 x i64] [i64 1027]
 561 #pragma omp target data map(close, tofrom \
 562                             : arg)
 563   {++arg;}
 564 }
 565 #endif
 566 ///==========================================================================///
 567 // RUN: %clang_cc1 -DCK7 -verify -Wno-vla  -fopenmp -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
 568 // RUN: %clang_cc1 -DCK7 -fopenmp -fopenmp-targets=x86_64 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 569 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK7 --check-prefix CK7-64
 570
 571 // RUN: %clang_cc1 -DCK7 -verify -Wno-vla  -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY7 %s
 572 // RUN: %clang_cc1 -DCK7 -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 573 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY7 %s
 574 // SIMD-ONLY7-NOT: {{__kmpc|__tgt}}
 575 #ifdef CK7
 576 // CK7: private unnamed_addr constant [2 x i64] [i64 64, i64 64]
 577 // CK7: private unnamed_addr constant [2 x i64] [i64 3, i64 64]
 578 // CK7-NOT: private unnamed_addr constant [2 x i64] [i64 64, i64 3]
 579 // CK7: test_device_ptr_addr
 580 void test_device_ptr_addr(int arg) {
 581   int *p;
 582   // CK7: add nsw i32
 583   // CK7: add nsw i32
 584   #pragma omp target data use_device_ptr(p) use_device_addr(arg)
 585   { ++arg, ++(*p); }
 586
 587   short x[10];
 588   short *xp = &x[0];
 589
 590   x[1] = 111;
 591
 592   #pragma omp target data map(tofrom: x) use_device_addr(xp[1:3])
 593   {
 594     xp[1] = 222;
 595   }
 596 }
 597 #endif
 598 ///==========================================================================///
 599 // RUN: %clang_cc1 -DCK8 -verify -Wno-vla  -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-64
 600 // RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 601 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-64
 602 // RUN: %clang_cc1 -DCK8 -verify -Wno-vla  -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-32
 603 // RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 604 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-32
 605
 606 // RUN: %clang_cc1 -DCK8 -verify -Wno-vla  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 607 // RUN: %clang_cc1 -DCK8 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 608 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 609 // RUN: %clang_cc1 -DCK8 -verify -Wno-vla  -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 610 // RUN: %clang_cc1 -DCK8 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 611 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 612 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK8
 613 #ifdef CK8
 614 struct S1 {
 615   int i;
 616 };
 617 struct S2 {
 618   S1 s;
 619   struct S2 *ps;
 620 };
 621
 622 void test_present_modifier(int arg) {
 623   S2 *ps1;
 624   S2 *ps2;
 625
 626   // Make sure the struct picks up present even if another element of the struct
 627   // doesn't have present.
 628   // CK8: private unnamed_addr constant [11 x i64] [i64 0, i64 {{4|8}}, i64 {{4|8}}, i64 4, i64 4, i64 4, i64 0, i64 4, i64 {{4|8}}, i64 {{4|8}}, i64 4]
 629   // CK8: private unnamed_addr constant [11 x i64]
 630
 631 // ps1
 632 //
 633 // PRESENT=0x1000 = 0x1000
 634 // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
 635 // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
 636 // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
 637 // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
 638 //
 639 // CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]],
 640 // CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]],
 641
 642 // arg
 643 //
 644 // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
 645 //
 646 // CK8-SAME: {{^}} i64 [[#0x1003]],
 647
 648 // ps2
 649 //
 650 // PRESENT=0x1000 = 0x1000
 651 // MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
 652 // MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
 653 // PTR_AND_OBJ=0x10 = 0x10
 654 // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
 655 //
 656 // CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
 657 // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
 658 #pragma omp target data map(tofrom         \
 659                             : ps1->s)      \
 660     map(present, tofrom                    \
 661         : arg, ps1->ps->ps->ps->s, ps2->s) \
 662         map(tofrom                         \
 663             : ps2->ps->ps->ps->s)
 664   {
 665     ++(arg);
 666   }
 667 }
 668 #endif
 669 ///==========================================================================///
 670 // RUN: %clang_cc1 -DCK9 -verify -Wno-vla  -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-64
 671 // RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 672 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-64
 673 // RUN: %clang_cc1 -DCK9 -verify -Wno-vla  -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
 674 // RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 675 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
 676
 677 // RUN: %clang_cc1 -DCK9 -verify -Wno-vla  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 678 // RUN: %clang_cc1 -DCK9 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 679 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 680 // RUN: %clang_cc1 -DCK9 -verify -Wno-vla  -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 681 // RUN: %clang_cc1 -DCK9 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 682 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla  %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 683 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 684 #ifdef CK9
 685 void test_present_modifier(int arg) {
 686 // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
 687 // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]]
 688 #pragma omp target data map(present, tofrom \
 689                             : arg)
 690   {++arg;}
 691 }
 692 #endif
 693 #endif