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