451 lines · cpp
1// expected-no-diagnostics2#ifndef HEADER3#define HEADER4 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-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 %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-649// 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-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 %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-3212 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 %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 %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s16// 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 %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 %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s19// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}20#ifdef CK121 22double *g;23 24// CK1: @g ={{.*}} global ptr25// 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: @_Z3foo39template<typename T>40void 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 146 // 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 nuw double, ptr [[TT]], i32 153 #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 nuw double, ptr [[TTT]], i32 160 ++g;61 62 // CK1: [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],63 // CK1: [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 064 // 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 nuw float, ptr [[TT1]], i32 171 #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 nuw float, ptr [[TTT]], i32 178 ++l;79 80 // CK1-NOT: call void @__tgt_target81 // CK1: [[TTT:%.+]] = load ptr, ptr [[DECL]],82 // CK1: getelementptr inbounds nuw float, ptr [[TTT]], i32 183 #pragma omp target data map(l[:10]) use_device_ptr(l) if(0)84 {85 ++l;86 }87 // CK1-NOT: call void @__tgt_target88 // CK1: [[TTT:%.+]] = load ptr, ptr [[DECL]],89 // CK1: getelementptr inbounds nuw float, ptr [[TTT]], i32 190 ++l;91 92 // CK1: [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],93 // CK1: [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 094 // 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 nuw float, ptr [[TT1]], i32 1101 #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 nuw float, ptr [[TTT]], i32 1108 ++l;109 110 // CK1: [[CMP:%.+]] = icmp ne ptr %{{.+}}, null111 // 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 0116 // 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 nuw float, ptr [[TT1]], i32 1123 // CK1: br label %[[BEND:.+]]124 125 // CK1: [[BELSE]]:126 // CK1: [[TTT:%.+]] = load ptr, ptr [[DECL]],127 // CK1: getelementptr inbounds nuw float, ptr [[TTT]], i32 1128 // 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 nuw float, ptr [[TTT]], i32 1146 ++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 0151 // 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 nuw float, ptr [[TT2]], i32 1160 #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 nuw float, ptr [[TTTT]], i32 1168 ++lr;169 170 // CK1: [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],171 // CK1: [[BP:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0172 // 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 nuw i32, ptr [[TT1]], i32 1179 #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 nuw i32, ptr [[TTT]], i32 1186 ++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 0191 // 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 nuw i32, ptr [[TT2]], i32 1200 #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 nuw i32, ptr [[TTTT]], i32 1208 ++tr;209 210 // CK1: [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],211 // CK1: [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 0212 // 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 nuw float, ptr [[TT1]], i32 1219 #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 nuw float, ptr [[TTT]], i32 1226 ++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 nuw float, ptr [[_TT1]], i32 1236 // CK1: [[TT1:%.+]] = load ptr, ptr [[PVT]],237 // CK1: getelementptr inbounds nuw i32, ptr [[TT1]], i32 1238 #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 nuw float, ptr [[_TTT]], i32 1245 // CK1: [[TTT:%.+]] = load ptr, ptr {{%.+}},246 // CK1: getelementptr inbounds nuw i32, ptr [[TTT]], i32 1247 ++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 nuw float, ptr [[_TT1]], i32 1256 // CK1: [[TT1:%.+]] = load ptr, ptr [[PVT]],257 // CK1: getelementptr inbounds nuw i32, ptr [[TT1]], i32 1258 #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 nuw float, ptr [[_TTT]], i32 1265 // CK1: [[TTT:%.+]] = load ptr, ptr {{%.+}},266 // CK1: getelementptr inbounds nuw i32, ptr [[TTT]], i32 1267 ++l; ++t;268 269 // CK1: [[T1:%.+]] = load ptr, ptr [[DECL:%.+]],270 // CK1: [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1271 // 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 nuw i32, ptr [[TT1]], i32 1278 #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 nuw i32, ptr [[TTT]], i32 1285 ++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 1290 // 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 nuw i32, ptr [[TT2]], i32 1299 #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 nuw i32, ptr [[TTTT]], i32 1307 ++l; ++tr;308 309}310 311void bar(float *&a, int *&b) {312 foo<int>(a,b);313}314 315#endif316///==========================================================================///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-64318// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s319// 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-64320// 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-32321// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s322// 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-32323 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 %s325// 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 %s326// 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 %s327// 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 %s328// 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 %s329// 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 %s330// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}331#ifdef CK2332 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 339template <typename T>340struct 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 1350 // 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 nuw double, ptr [[TT2]], i32 1358 #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 nuw [[ST]], ptr %this1, i32 0, i32 0364 // CK2: [[TTT:%.+]] = load ptr, ptr [[DECL]],365 // CK2: getelementptr inbounds nuw double, ptr [[TTT]], i32 1366 a++;367 368 // CK2: [[BP:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 1369 // 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 nuw double, ptr [[TT2]], i32 1377 #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 nuw [[ST]], ptr %{{.+}}, i32 0, i32 1383 // CK2: [[TTT:%.+]] = load ptr, ptr [[DECL]],384 // CK2: [[TTTT:%.+]] = load ptr, ptr [[TTT]],385 // CK2: getelementptr inbounds nuw double, ptr [[TTTT]], i32 1386 b++;387 388 // CK2: [[BP:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2389 // 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 nuw double, ptr [[TT2]], i32 1397 #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 nuw [[ST]], ptr %this1, i32 0, i32 0404 // CK2: [[TTT:%.+]] = load ptr, ptr [[DECL]],405 // CK2: getelementptr inbounds nuw double, ptr [[TTT]], i32 1406 a++;407 la++;408 409 // CK2: [[BP1:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 1410 // CK2: store ptr [[RVAL1:%.+]], ptr [[BP1]],411 // CK2: [[BP2:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2412 // 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 nuw double, ptr [[_TT2]], i32 1423 // CK2: [[TT1:%.+]] = load ptr, ptr [[_PVT1]],424 // CK2: [[_TT1:%.+]] = load ptr, ptr [[TT1]],425 // CK2: getelementptr inbounds nuw double, ptr [[_TT1]], i32 1426 #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 nuw [[ST]], ptr %this1, i32 0, i32 0433 // CK2: [[TTT:%.+]] = load ptr, ptr [[DECL]],434 // CK2: getelementptr inbounds nuw double, ptr [[TTT]], i32 1435 // CK2: [[_DECL:%.+]] = getelementptr inbounds nuw [[ST]], ptr %this1, i32 0, i32 1436 // CK2: [[_TTT:%.+]] = load ptr, ptr [[_DECL]],437 // CK2: [[_TTTT:%.+]] = load ptr, ptr [[_TTT]],438 // CK2: getelementptr inbounds nuw double, ptr [[_TTTT]], i32 1439 a++;440 b++;441 }442};443 444void bar(double *arg){445 ST<double> A(arg);446 A.foo(arg);447 ++arg;448}449#endif450#endif451