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