694 lines · cpp
1// expected-no-diagnostics2#ifndef HEADER3#define HEADER4 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-647// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s8// 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-649// 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-3210// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s11// 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-3212 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 %s14// 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 %s15// 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 %s16// 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 %s17// 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 %s18// 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 %s19// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}20#ifdef CK121 22// CK1: [[ST:%.+]] = type { i32, ptr }23template <typename T>24struct ST {25 T a;26 double *b;27};28 29ST<int> gb;30double 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: _Z3fooi48void foo(int arg) {49 int la;50 float lb[arg];51 52 // Region 0053 // 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 i6455 // 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{{.+}} 060 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 061 // CK1-DAG: store ptr @gc, ptr [[BP0]]62 // CK1-DAG: store ptr @gc, ptr [[P0]]63 64 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 165 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 0173 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 174 #pragma omp target data map(la) if(1+3-4)75 {++arg;}76 77 // Region 0278 // 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{{.+}} 085 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 086 // 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 %{{[^,]+}}, 194 // 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 03108 // 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{{.+}} 0114 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0115 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0116 // 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 %{{[^,]+}}, 4120 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64121 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4122 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1123 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 %{{[^,]+}}, 1132 {++arg;}133 134 // Region 04135 // 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{{.+}} 0141 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0142 // CK1-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PSZ]], i{{.+}} 0, i{{.+}} 0143 // 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{{.+}} 1149 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1150 // 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{{.+}} 0153 // CK1-DAG: [[SEC11]] = load ptr, ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1),154 155 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1156 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 %{{[^,]+}}, 1165 {++arg;}166 167 // Region 05168 // 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{{.+}} 0174 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0175 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0176 // 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 %{{[^,]+}}, 4180 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64181 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4182 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1183 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 %{{[^,]+}}, 1192 {++arg;}193 194 // Region 06195 // 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{{.+}} 0201 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0202 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0203 // 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 %{{[^,]+}}, 4207 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64208 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4209 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1210 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#endif220///==========================================================================///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-64222// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s223// 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-64224// 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-32225// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s226// 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-32227 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 %s229// 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 %s230// 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 %s231// 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 %s232// 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 %s233// 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 %s234// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}235#ifdef CK1A236 237// CK1A: [[ST:%.+]] = type { i32, ptr }238template <typename T>239struct ST {240 T a;241 double *b;242};243 244ST<int> gb;245double gc[100];246 247// PRESENT=0x1000 | TO=0x1 = 0x1001248// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]249 250// TO=0x1 = 0x1251// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1]]]252 253// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405254// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]255 256// CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x405257// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x405]]]258 259// CK1A-LABEL: _Z3fooi260void foo(int arg) {261 int la;262 float lb[arg];263 264 // Region 00265 // 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{{.+}} 0271 // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0272 // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0273 // 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 %{{[^,]+}}, 4277 // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64278 // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4279 // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1280 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 01289 // 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{{.+}} 0295 // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0296 // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0297 // 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 %{{[^,]+}}, 4301 // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64302 // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4303 // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1304 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#endif314///==========================================================================///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-64316// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s317// 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-64318// 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-32319// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s320// 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-32321 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 %s323// 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 %s324// 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 %s325// 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 %s326// 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 %s327// 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 %s328// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}329#ifdef CK2330 331// CK2: [[ST:%.+]] = type { i32, ptr }332template <typename T>333struct ST {334 T a;335 double *b;336 337 T foo(T arg) {338 // Region 00339 #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: _Z3bari349int bar(int arg){350 ST<int> A;351 return A.foo(arg);352}353 354// Region 00355// CK2-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64356// 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{{.+}} 0365// CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0366// CK2-DAG: [[PS0:%.+]] = getelementptr inbounds [2 x i64], ptr [[PS]], i{{.+}} 0, i{{.+}} 0367// 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 1371 372// CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1373// CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1374// CK2-DAG: store ptr [[SEC0]], ptr [[BP1]]375// CK2-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]376// CK2-DAG: [[SEC1]] = getelementptr inbounds {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 1377// CK2-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],378// CK2-DAG: [[SEC111]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1379 380// CK2: br label %[[IFEND:[^,]+]]381 382// CK2: [[IFELSE]]383// CK2: br label %[[IFEND]]384// CK2: [[IFEND]]385// CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1386// 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#endif398///==========================================================================///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-64400// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s401// 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-64402// 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-32403// RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s404// 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-32405 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 %s407// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s408// 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 %s409// RUN: %clang_cc1 -DCK3 -verify -Wno-vla -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s410// RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s411// 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 %s412// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}413#ifdef CK3414 415// CK3-LABEL: no_target_devices416void no_target_devices(int arg) {417 // CK3-NOT: tgt_target_data_begin418 // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1419 // CK3-NOT: tgt_target_data_end420 // CK3: ret421 #pragma omp target data map(to: arg) if(arg) device(4)422 {++arg;}423}424#endif425///==========================================================================///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-64427// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s428// 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-64429// 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-32430// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s431// 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-32432 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 %s434// 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 %s435// 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 %s436// 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 %s437// 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 %s438// 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 %s439// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}440#ifdef CK4441 442// CK4: [[STT:%.+]] = type { i32, ptr }443template <typename T>444struct STT {445 T a;446 double *b;447 448 T foo(T arg) {449 // Region 00450 #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: _Z3bari460int bar(int arg){461 STT<int> A;462 return A.foo(arg);463}464 465// Region 00466// CK4-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64467// 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{{.+}} 0476// CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0477// CK4-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i{{.+}} 0, i{{.+}} 0478// 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 1482 483// CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1484// CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1485// CK4-DAG: store ptr [[SEC0]], ptr [[BP1]]486// CK4-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]487// CK4-DAG: [[SEC1]] = getelementptr inbounds {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 1488// CK4-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],489// CK4-DAG: [[SEC111]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1490 491// CK4: br label %[[IFEND:[^,]+]]492 493// CK4: [[IFELSE]]494// CK4: br label %[[IFEND]]495// CK4: [[IFEND]]496// CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1497// 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#endif509///==========================================================================///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-64511// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s512// 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-64513// 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-32514// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s515// 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-32516 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 %s518// RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s519// 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 %s520// RUN: %clang_cc1 -DCK5 -verify -Wno-vla -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s521// RUN: %clang_cc1 -DCK5 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s522// 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 %s523// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK5524#ifdef CK5525struct S1 {526 int i;527};528struct S2 {529 S1 s;530 struct S2 *ps;531};532 533void 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#endif543///==========================================================================///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-64545// RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s546// 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-64547// 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-32548// RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s549// 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-32550 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 %s552// RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s553// 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 %s554// RUN: %clang_cc1 -DCK6 -verify -Wno-vla -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s555// RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s556// 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 %s557// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}558#ifdef CK6559void 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#endif566///==========================================================================///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-64568// RUN: %clang_cc1 -DCK7 -fopenmp -fopenmp-targets=x86_64 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s569// 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-64570 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 %s572// RUN: %clang_cc1 -DCK7 -fopenmp-simd -fopenmp-targets=x86_64 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s573// 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 %s574// SIMD-ONLY7-NOT: {{__kmpc|__tgt}}575#ifdef CK7576// 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_addr580void test_device_ptr_addr(int arg) {581 int *p;582 // CK7: add nsw i32583 // CK7: add nsw i32584 #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#endif598///==========================================================================///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-64600// RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s601// 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-64602// 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-32603// RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s604// 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-32605 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 %s607// RUN: %clang_cc1 -DCK8 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s608// 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 %s609// RUN: %clang_cc1 -DCK8 -verify -Wno-vla -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s610// RUN: %clang_cc1 -DCK8 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s611// 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 %s612// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK8613#ifdef CK8614struct S1 {615 int i;616};617struct S2 {618 S1 s;619 struct S2 *ps;620};621 622void 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 struct627 // 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// ps1632//633// PRESENT=0x1000 = 0x1000634// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010635// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010636// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013637// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003638//639// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]],640// CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]],641 642// arg643//644// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003645//646// CK8-SAME: {{^}} i64 [[#0x1003]],647 648// ps2649//650// PRESENT=0x1000 = 0x1000651// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003652// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010653// PTR_AND_OBJ=0x10 = 0x10654// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13655//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#endif669///==========================================================================///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-64671// RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s672// 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-64673// 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-32674// RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s675// 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-32676 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 %s678// RUN: %clang_cc1 -DCK9 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s679// 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 %s680// RUN: %clang_cc1 -DCK9 -verify -Wno-vla -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s681// RUN: %clang_cc1 -DCK9 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s682// 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 %s683// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}684#ifdef CK9685void test_present_modifier(int arg) {686// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003687// CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]]688#pragma omp target data map(present, tofrom \689 : arg)690 {++arg;}691}692#endif693#endif694