diff options
| author | Samuel Antao <sfantao@us.ibm.com> | 2016-04-27 22:40:57 +0000 |
|---|---|---|
| committer | Samuel Antao <sfantao@us.ibm.com> | 2016-04-27 22:40:57 +0000 |
| commit | d62c9cc77260e3459d4788394d379c0e1e04963d (patch) | |
| tree | bc0b8f361c3271a5ebd0eb69c4163dbc17f2eaa0 /test/OpenMP/target_map_codegen.cpp | |
| parent | e5c389d29dfe747736ec899b3885b1f49b0d71b1 (diff) | |
[OpenMP] Map clause codegeneration.
Summary:
Implement codegen for the map clause. All the new list items in 4.5 specification are supported.
Fix bug in the generation of array sections that was exposed by some of the map clause tests: for pointer types the offsets have to be calculated from the pointee not the pointer.
Reviewers: hfinkel, kkwli0, carlo.bertolli, arpith-jacob, ABataev
Subscribers: ABataev, cfe-commits, caomhin, fraggamuffin
Differential Revision: http://reviews.llvm.org/D16749
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@267808 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/OpenMP/target_map_codegen.cpp')
| -rw-r--r-- | test/OpenMP/target_map_codegen.cpp | 3252 |
1 files changed, 3250 insertions, 2 deletions
diff --git a/test/OpenMP/target_map_codegen.cpp b/test/OpenMP/target_map_codegen.cpp index 3a59c8f522..03a514d11f 100644 --- a/test/OpenMP/target_map_codegen.cpp +++ b/test/OpenMP/target_map_codegen.cpp @@ -454,8 +454,8 @@ void implicit_maps_array (int a){ #ifdef CK10 // CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] -// Map types: OMP_MAP_BYCOPY | OMP_MAP_PTR = 128 + 32 -// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 160] +// Map types: OMP_MAP_BYCOPY = 128 +// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] // CK10-LABEL: implicit_maps_pointer void implicit_maps_pointer (){ @@ -1012,4 +1012,3252 @@ void implicit_maps_template_type_capture (int a){ // CK18-32: {{.+}} = load i32, i32* [[ADDR]], #endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32 +// RUN: %clang_cc1 -DCK19 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK19 --check-prefix CK19-32 +#ifdef CK19 + +// CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1] + +// CK19: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240] +// CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 2] + +// CK19: [[SIZE18:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240] +// CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3] + +// CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 0] + +// CK19: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1] + +// CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3] + +// CK19: [[SIZE22:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3] + +// CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 7] + +// CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 480] +// CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 24] +// CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 16] +// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99] + +// CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4] +// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99] + +// CK19: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i32] [i32 128, i32 128, i32 128, i32 3] + +// CK19: [[SIZE31:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 40] +// CK19: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i32] [i32 128, i32 128, i32 128, i32 3] + +// CK19: [[SIZE32:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728] +// CK19: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE33:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728] +// CK19: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE34:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728] +// CK19: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE36:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 208] +// CK19: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3] + +// CK19: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3] + +// CK19: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3] + +// CK19: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3] + +// CK19: [[SIZE41:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 208] +// CK19: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3] + +// CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 104] +// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99] + +// CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19-LABEL: explicit_maps_single +void explicit_maps_single (int ii){ + // Map of a scalar. + int a = ii; + + // Region 00 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK19: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:a) + { + ++a; + } + + // Map of an array. + int arra[100]; + + // Region 01 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [100 x i32]* [[VAR0]] to i8* + + // CK19: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(to:arra) + { + arra[50]++; + } + + // Region 02 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 20 + + // CK19: call void [[CALL02:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(from:arra[20:60]) + { + arra[50]++; + } + + // Region 03 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(tofrom:arra[:60]) + { + arra[50]++; + } + + // Region 04 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL04:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(alloc:arra[:]) + { + arra[50]++; + } + + // Region 05 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 15 + + // CK19: call void [[CALL05:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(to:arra[15]) + { + arra[15]++; + } + + // Region 06 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}} + + // CK19: call void [[CALL06:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(tofrom:arra[ii:ii+23]) + { + arra[50]++; + } + + // Region 07 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL07:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(alloc:arra[:ii]) + { + arra[50]++; + } + + // Region 08 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}} + + // CK19: call void [[CALL08:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(tofrom:arra[ii]) + { + arra[15]++; + } + + // Map of a pointer. + int *pa; + + // Region 09 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32** [[VAR0]] to i8* + + // CK19: call void [[CALL09:@.+]](i32** {{[^,]+}}) + #pragma omp target map(from:pa) + { + pa[50]++; + } + + // Region 10 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 20 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL10:@.+]](i32* {{[^,]+}}) + #pragma omp target map(tofrom:pa[20:60]) + { + pa[50]++; + } + + // Region 11 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL11:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:pa[:60]) + { + pa[50]++; + } + + // Region 12 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 15 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL12:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:pa[15]) + { + pa[15]++; + } + + // Region 13 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL13:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:pa[ii-23:ii]) + { + pa[50]++; + } + + // Region 14 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL14:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:pa[:ii]) + { + pa[50]++; + } + + // Region 15 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL15:@.+]](i32* {{[^,]+}}) + #pragma omp target map(from:pa[ii+12]) + { + pa[15]++; + } + + // Map of a variable-size array. + int va[ii]; + + // Region 16 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE16]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[VAR1]] to i8* + // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4 + + // CK19: call void [[CALL16:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(to:va) + { + va[50]++; + } + + // Region 17 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE17]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE17]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 20 + + // CK19: call void [[CALL17:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(from:va[20:60]) + { + va[50]++; + } + + // Region 18 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE18]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE18]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0 + + // CK19: call void [[CALL18:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(tofrom:va[:60]) + { + va[50]++; + } + + // Region 19 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE19]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0 + + // CK19: call void [[CALL19:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(alloc:va[:]) + { + va[50]++; + } + + // Region 20 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE20]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE20]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 15 + + // CK19: call void [[CALL20:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(to:va[15]) + { + va[15]++; + } + + // Region 21 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE21]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}} + + // CK19: call void [[CALL21:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(tofrom:va[ii:ii+23]) + { + va[50]++; + } + + // Region 22 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE22]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE22]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}} + + // CK19: call void [[CALL22:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(tofrom:va[ii]) + { + va[15]++; + } + + // Always. + // Region 23 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE23]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE23]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK19: call void [[CALL23:@.+]](i32* {{[^,]+}}) + #pragma omp target map(always, tofrom: a) + { + a++; + } + + // Multidimensional arrays. + int marr[4][5][6]; + int ***mptr; + + // Region 24 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE24]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE24]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0]] to i8* + + // CK19: call void [[CALL24:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr) + { + marr[1][2][3]++; + } + + // Region 25 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE25]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE25]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK19: call void [[CALL25:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr[1][2][2:4]) + { + marr[1][2][3]++; + } + + // Region 26 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE26]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE26]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK19: call void [[CALL26:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr[1][2][:]) + { + marr[1][2][3]++; + } + + // Region 27 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE27]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE27]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK19: call void [[CALL27:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr[1][2][3]) + { + marr[1][2][3]++; + } + + // Region 28 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE28]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE28]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32*** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32*** [[SEC0:%.+]] to i8* + // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32*** [[SEC0]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32** [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32** [[SEC11:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC11]] = load i32**, i32*** [[SEC111:%[^,]+]], + // CK19-DAG: [[SEC111]] = getelementptr {{.*}}i32*** [[SEC1111:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC1111]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = bitcast i32** [[SEC1]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast i32* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.*}}i32* [[SEC22:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC22]] = load i32*, i32** [[SEC222:%[^,]+]], + // CK19-DAG: [[SEC222]] = getelementptr {{.*}}i32** [[SEC2222:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC2222]] = load i32**, i32*** [[SEC22222:%[^,]+]], + // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}i32*** [[SEC222222:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC222222]] = load i32***, i32**** [[PTR]], + + // CK19: call void [[CALL28:@.+]](i32*** {{[^,]+}}) + #pragma omp target map(tofrom: mptr[1][2][2:4]) + { + mptr[1][2][3]++; + } + + // Region 29 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE29]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE29]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32*** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32*** [[SEC0:%.+]] to i8* + // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32*** [[SEC0]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32** [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32** [[SEC11:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC11]] = load i32**, i32*** [[SEC111:%[^,]+]], + // CK19-DAG: [[SEC111]] = getelementptr {{.*}}i32*** [[SEC1111:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC1111]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = bitcast i32** [[SEC1]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast i32* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.*}}i32* [[SEC22:[^,]+]], i{{.+}} 3 + // CK19-DAG: [[SEC22]] = load i32*, i32** [[SEC222:%[^,]+]], + // CK19-DAG: [[SEC222]] = getelementptr {{.*}}i32** [[SEC2222:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC2222]] = load i32**, i32*** [[SEC22222:%[^,]+]], + // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}i32*** [[SEC222222:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC222222]] = load i32***, i32**** [[PTR]], + + // CK19: call void [[CALL29:@.+]](i32*** {{[^,]+}}) + #pragma omp target map(tofrom: mptr[1][2][3]) + { + mptr[1][2][3]++; + } + + // Multidimensional VLA. + double mva[23][ii][ii+5]; + + // Region 30 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE30]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 23 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 23 to i8*), i8** [[P0]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // CK19-64-DAG: [[VAR1]] = zext i32 %{{[^,]+}} to i64 + // CK19-64-DAG: [[VAR11]] = zext i32 %{{[^,]+}} to i64 + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S2]] + // CK19-DAG: [[CBPVAL2]] = inttoptr i[[Z]] [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = inttoptr i[[Z]] [[VAR22:%.+]] to i8* + // CK19-64-DAG: [[VAR2]] = zext i32 %{{[^,]+}} to i64 + // CK19-64-DAG: [[VAR22]] = zext i32 %{{[^,]+}} to i64 + // + // CK19-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]] + // CK19-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]] + // CK19-DAG: store i[[Z]] [[CSVAL3:%[^,]+]], i[[Z]]* [[S3]] + // CK19-DAG: [[CBPVAL3]] = bitcast double* [[VAR3:%.+]] to i8* + // CK19-DAG: [[CPVAL3]] = bitcast double* [[VAR3]] to i8* + // CK19-DAG: [[CSVAL3]] = mul nuw i[[Z]] %{{[^,]+}}, {{8|4}} + + // CK19: call void [[CALL30:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, double* %{{[^,]+}}) + #pragma omp target map(tofrom: mva) + { + mva[1][2][3]++; + } + + // Region 31 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE31]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE31]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 23 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 23 to i8*), i8** [[P0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = inttoptr i[[Z]] [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = inttoptr i[[Z]] [[VAR22:%.+]] to i8* + // + // CK19-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]] + // CK19-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]] + // CK19-DAG: [[CBPVAL3]] = bitcast double* [[VAR3:%.+]] to i8* + // CK19-DAG: [[CPVAL3]] = bitcast double* [[SEC3:%.+]] to i8* + // CK19-DAG: [[SEC3]] = getelementptr {{.*}}double* [[SEC33:%.+]], i[[Z]] 0 + // CK19-DAG: [[SEC33]] = getelementptr {{.*}}double* [[SEC333:%.+]], i[[Z]] [[IDX3:%.+]] + // CK19-DAG: [[IDX3]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}} + // CK19-DAG: [[SEC333]] = getelementptr {{.*}}double* [[VAR3]], i[[Z]] [[IDX33:%.+]] + // CK19-DAG: [[IDX33]] = mul nsw i[[Z]] 1, %{{[^,]+}} + + // CK19: call void [[CALL31:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, double* %{{[^,]+}}) + #pragma omp target map(tofrom: mva[1][ii-2][:5]) + { + mva[1][2][3]++; + } + + // Multidimensional array sections. + double marras[11][12][13]; + double mvlaas[11][ii][13]; + double ***mptras; + + // Region 32 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE32]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE32]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0]] to i8* + + // CK19: call void [[CALL32:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) + #pragma omp target map(marras) + { + marras[1][2][3]++; + } + + // Region 33 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE33]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE33]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [12 x [13 x double]]* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 0 + + // CK19: call void [[CALL33:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) + #pragma omp target map(marras[:]) + { + marras[1][2][3]++; + } + + // Region 34 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE34]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE34]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [12 x [13 x double]]* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 0 + + // CK19: call void [[CALL34:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) + #pragma omp target map(marras[:][:][:]) + { + marras[1][2][3]++; + } + + // Region 35 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE35]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i[[Z]] [[CSVAL0:%[^,]+]], i[[Z]]* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [13 x double]* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[12 x [13 x double]]* [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 0 + // CK19-DAG: [[SEC00]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 1 + // CK19-DAG: [[CSVAL0]] = mul nuw i[[Z]] %{{[^,]+}}, 104 + + // CK19: call void [[CALL35:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) + #pragma omp target map(marras[1][:ii][:]) + { + marras[1][2][3]++; + } + + // Region 36 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE36]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE36]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [13 x double]* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[13 x double]* [[SEC00:%[^,]+]], i{{.+}} 0 + // CK19-DAG: [[SEC00]] = getelementptr {{.+}}[12 x [13 x double]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[SEC000]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL36:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) + #pragma omp target map(marras[:1][:2][:13]) + { + marras[1][2][3]++; + } + + // Region 37 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE37]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[P0]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: store i[[Z]] [[CSVAL2:%[^,]+]], i[[Z]]* [[S2]] + // CK19-DAG: [[CBPVAL2]] = bitcast [13 x double]* [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast [13 x double]* [[VAR2]] to i8* + // CK19-DAG: [[CSVAL2]] = mul nuw i[[Z]] %{{[^,]+}}, 104 + + // CK19: call void [[CALL37:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) + #pragma omp target map(mvlaas) + { + mvlaas[1][2][3]++; + } + + // Region 38 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE38]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[P0]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: store i[[Z]] [[CSVAL2:%[^,]+]], i[[Z]]* [[S2]] + // CK19-DAG: [[CBPVAL2]] = bitcast [13 x double]* [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast [13 x double]* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC22:%[^,]+]] + // CK19-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}} + // CK19-DAG: [[CSVAL2]] = mul nuw i[[Z]] %{{[^,]+}}, 104 + + // CK19: call void [[CALL38:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) + #pragma omp target map(mvlaas[:]) + { + mvlaas[1][2][3]++; + } + + // Region 39 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE39]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[P0]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: store i[[Z]] [[CSVAL2:%[^,]+]], i[[Z]]* [[S2]] + // CK19-DAG: [[CBPVAL2]] = bitcast [13 x double]* [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast [13 x double]* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC22:%[^,]+]] + // CK19-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}} + // CK19-DAG: [[CSVAL2]] = mul nuw i[[Z]] %{{[^,]+}}, 104 + + // CK19: call void [[CALL39:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) + #pragma omp target map(mvlaas[:][:][:]) + { + mvlaas[1][2][3]++; + } + + // Region 40 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE40]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[P0]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i[[Z]] {{8|4}}, i[[Z]]* [[S1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: store i[[Z]] [[CSVAL2:%[^,]+]], i[[Z]]* [[S2]] + // CK19-DAG: [[CBPVAL2]] = bitcast [13 x double]* [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast [13 x double]* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[SEC22:%[^,]+]], i[[Z]] 0 + // CK19-DAG: [[SEC22]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC222:%[^,]+]] + // CK19-DAG: [[SEC222]] = mul nsw i[[Z]] 1, %{{[^,]+}} + + // CK19: call void [[CALL40:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) + #pragma omp target map(mvlaas[1][:ii][:]) + { + mvlaas[1][2][3]++; + } + + // Region 41 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE41]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE41]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[BP0]] + // CK19-DAG: store i8* inttoptr (i[[Z]] 11 to i8*), i8** [[P0]] + // + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = inttoptr i[[Z]] [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = inttoptr i[[Z]] [[VAR11:%.+]] to i8* + // + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = bitcast [13 x double]* [[VAR2:%.+]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast [13 x double]* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[SEC22:%[^,]+]], i[[Z]] 0 + // CK19-DAG: [[SEC22]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC222:%[^,]+]] + // CK19-DAG: [[SEC222]] = mul nsw i[[Z]] 0, %{{[^,]+}} + + // CK19: call void [[CALL41:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) + #pragma omp target map(mvlaas[:1][:2][:13]) + { + mvlaas[1][2][3]++; + } + + // Region 42 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE42]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE42]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast double*** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast double*** [[SEC0:%.+]] to i8* + // CK19-DAG: [[VAR0]] = load double***, double**** [[PTR:%[^,]+]], + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}double*** [[SEC00:[^,]+]], i{{.+}} 0 + // CK19-DAG: [[SEC00]] = load double***, double**** [[PTR]], + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast double*** [[SEC0]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast double** [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}double** [[SEC11:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC11]] = load double**, double*** [[SEC111:%[^,]+]], + // CK19-DAG: [[SEC111]] = getelementptr {{.*}}double*** [[SEC1111:[^,]+]], i{{.+}} 0 + // CK19-DAG: [[SEC1111]] = load double***, double**** [[PTR]], + + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = bitcast double** [[SEC1]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast double* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.*}}double* [[SEC22:[^,]+]], i{{.+}} 0 + // CK19-DAG: [[SEC22]] = load double*, double** [[SEC222:%[^,]+]], + // CK19-DAG: [[SEC222]] = getelementptr {{.*}}double** [[SEC2222:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC2222]] = load double**, double*** [[SEC22222:%[^,]+]], + // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}double*** [[SEC222222:[^,]+]], i{{.+}} 0 + // CK19-DAG: [[SEC222222]] = load double***, double**** [[PTR]], + + // CK19: call void [[CALL42:@.+]](double*** {{[^,]+}}) + #pragma omp target map(mptras[:1][2][:13]) + { + mptras[1][2][3]++; + } + + // Region 43 - the memory is not contiguous for this map - will map the whole last dimension. + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE43]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + // + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i[[Z]] [[CSVAL0:%[^,]+]], i[[Z]]* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [11 x [12 x [13 x double]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [13 x double]* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[12 x [13 x double]]* [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 0 + // CK19-DAG: [[SEC00]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 1 + // CK19-DAG: [[CSVAL0]] = mul nuw i[[Z]] %{{[^,]+}}, 104 + + // CK19: call void [[CALL43:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) + #pragma omp target map(marras[1][:ii][1:]) + { + marras[1][2][3]++; + } + +} + +// CK19: define {{.+}}[[CALL00]] +// CK19: define {{.+}}[[CALL01]] +// CK19: define {{.+}}[[CALL02]] +// CK19: define {{.+}}[[CALL03]] +// CK19: define {{.+}}[[CALL04]] +// CK19: define {{.+}}[[CALL05]] +// CK19: define {{.+}}[[CALL06]] +// CK19: define {{.+}}[[CALL07]] +// CK19: define {{.+}}[[CALL08]] +// CK19: define {{.+}}[[CALL09]] +// CK19: define {{.+}}[[CALL10]] +// CK19: define {{.+}}[[CALL11]] +// CK19: define {{.+}}[[CALL12]] +// CK19: define {{.+}}[[CALL13]] +// CK19: define {{.+}}[[CALL14]] +// CK19: define {{.+}}[[CALL15]] +// CK19: define {{.+}}[[CALL16]] +// CK19: define {{.+}}[[CALL17]] +// CK19: define {{.+}}[[CALL18]] +// CK19: define {{.+}}[[CALL19]] +// CK19: define {{.+}}[[CALL20]] +// CK19: define {{.+}}[[CALL21]] +// CK19: define {{.+}}[[CALL22]] +// CK19: define {{.+}}[[CALL23]] +// CK19: define {{.+}}[[CALL24]] +// CK19: define {{.+}}[[CALL25]] +// CK19: define {{.+}}[[CALL26]] +// CK19: define {{.+}}[[CALL27]] +// CK19: define {{.+}}[[CALL28]] +// CK19: define {{.+}}[[CALL29]] +// CK19: define {{.+}}[[CALL30]] +// CK19: define {{.+}}[[CALL31]] +// CK19: define {{.+}}[[CALL32]] +// CK19: define {{.+}}[[CALL33]] +// CK19: define {{.+}}[[CALL34]] +// CK19: define {{.+}}[[CALL35]] +// CK19: define {{.+}}[[CALL36]] +// CK19: define {{.+}}[[CALL37]] +// CK19: define {{.+}}[[CALL38]] +// CK19: define {{.+}}[[CALL39]] +// CK19: define {{.+}}[[CALL40]] +// CK19: define {{.+}}[[CALL41]] +// CK19: define {{.+}}[[CALL42]] +// CK19: define {{.+}}[[CALL43]] + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64 +// RUN: %clang_cc1 -DCK20 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK20 --check-prefix CK20-64 +// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-32 +// RUN: %clang_cc1 -DCK20 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK20 --check-prefix CK20-32 +#ifdef CK20 + +// CK20: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK20: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 12] +// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK20-LABEL: explicit_maps_references_and_function_args +void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){ + + int &aa = a; + float &bb = b; + int (&cc)[10] = c; + float *&dd = d; + + // Region 00 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast i32* [[RVAR00:%.+]] to i8* + // CK20-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK20-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK20: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:aa) + { + aa += 1; + } + + // Region 01 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast [10 x i32]* [[RVAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK20-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[RVAR00:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[RVAR0]] = load [10 x i32]*, [10 x i32]** [[VAR0:%[^,]+]] + // CK20-DAG: [[RVAR00]] = load [10 x i32]*, [10 x i32]** [[VAR0]] + + // CK20: call void [[CALL01:@.+]]([10 x i32]* {{[^,]+}}) + #pragma omp target map(to:cc[:5]) + { + cc[3] += 1; + } + + // Region 02 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast float* [[VAR0]] to i8* + + // CK20: call void [[CALL02:@.+]](float* {{[^,]+}}) + #pragma omp target map(from:b) + { + b += 1.0f; + } + + // Region 03 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast float* [[RVAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8* + // CK20-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]] + // CK20-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2 + // CK20-DAG: [[RVAR00]] = load float*, float** [[VAR0]] + + // CK20: call void [[CALL03:@.+]](float* {{[^,]+}}) + #pragma omp target map(from:d[2:3]) + { + d[2] += 1.0f; + } +} + +// CK20: define {{.+}}[[CALL00]] +// CK20: define {{.+}}[[CALL01]] +// CK20: define {{.+}}[[CALL02]] +// CK20: define {{.+}}[[CALL03]] + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64 +// RUN: %clang_cc1 -DCK21 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK21 --check-prefix CK21-64 +// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-32 +// RUN: %clang_cc1 -DCK21 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK21 --check-prefix CK21-32 +#ifdef CK21 +// CK21: [[ST:%.+]] = type { i32, i32, float* } + +// CK21: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK21: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492] +// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK21: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 500] +// CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 2, i32 98] + +// CK21: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492] +// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK21: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK21: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] 4, i[[Z]] 4] +// CK21: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 67] + +// CK21-LABEL: explicit_maps_template_args_and_members + +template <int X, typename T> +struct CC { + T A; + int A2; + float *B; + + int foo(T arg) { + float la[X]; + T *lb; + + // Region 00 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0:%.+]], i{{.+}} 0, i{{.+}} 0 + + // CK21: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(A) + { + A += 1; + } + + // Region 01 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK21-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 + // CK21-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK21: call void [[CALL01:@.+]](i32* {{[^,]+}}) + #pragma omp target map(lb[:X]) + { + lb[4] += 1; + } + + // Region 02 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast float** [[SEC0:%.+]] to i8* + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + + // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK21-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK21-DAG: [[CBPVAL1]] = bitcast float** [[SEC0]] to i8* + // CK21-DAG: [[CPVAL1]] = bitcast float* [[SEC1:%.+]] to i8* + // CK21-DAG: [[SEC1]] = getelementptr {{.*}}float* [[RVAR1:%[^,]+]], i{{.+}} 123 + // CK21-DAG: [[RVAR1]] = load float*, float** [[SEC1_:%[^,]+]] + // CK21-DAG: [[SEC1_]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + + // CK21: call void [[CALL02:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(from:B[X:X+2]) + { + B[2] += 1.0f; + } + + // Region 03 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [123 x float]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast [123 x float]* [[VAR0]] to i8* + + // CK21: call void [[CALL03:@.+]]([123 x float]* {{[^,]+}}) + #pragma omp target map(from:la) + { + la[3] += 1.0f; + } + + // Region 04 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK21: call void [[CALL04:@.+]](i32* {{[^,]+}}) + #pragma omp target map(from:arg) + { + arg +=1; + } + + // Make sure the extra flag is passed to the second map. + // Region 05 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK21-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK21-DAG: [[CBPVAL1]] = bitcast [[ST]]* [[VAR1:%.+]] to i8* + // CK21-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK21-DAG: [[SEC1]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK21: call void [[CALL05:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(A, A2) + { + A += 1; + A2 += 1; + } + return A; + } +}; + +int explicit_maps_template_args_and_members(int a){ + CC<123,int> c; + return c.foo(a); +} + +// CK21: define {{.+}}[[CALL00]] +// CK21: define {{.+}}[[CALL01]] +// CK21: define {{.+}}[[CALL02]] +// CK21: define {{.+}}[[CALL03]] +// CK21: define {{.+}}[[CALL04]] +// CK21: define {{.+}}[[CALL05]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64 +// RUN: %clang_cc1 -DCK22 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK22 --check-prefix CK22-64 +// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-32 +// RUN: %clang_cc1 -DCK22 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK22 --check-prefix CK22-32 +#ifdef CK22 + +// CK22-DAG: [[ST:%.+]] = type { float } +// CK22-DAG: [[STT:%.+]] = type { i32 } + +// CK22: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +int a; +int c[100]; +int *d; + +struct ST { + float fa; +}; + +ST sa ; +ST sc[100]; +ST *sd; + +template<typename T> +struct STT { + T fa; +}; + +STT<int> sta ; +STT<int> stc[100]; +STT<int> *std; + +// CK22-LABEL: explicit_maps_globals +int explicit_maps_globals(void){ + // Region 00 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast (i32* @a to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast (i32* @a to i8*), i8** [[P0]] + + // CK22: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(a) + { a+=1; } + + // Region 01 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[P0]] + + // CK22: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(c) + { c[3]+=1; } + + // Region 02 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast (i32** @d to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast (i32** @d to i8*), i8** [[P0]] + + // CK22: call void [[CALL02:@.+]](i32** {{[^,]+}}) + #pragma omp target map(d) + { d[3]+=1; } + + // Region 03 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast (i32* getelementptr inbounds ([100 x i32], [100 x i32]* @c, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]] + + // CK22: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(c[1:4]) + { c[3]+=1; } + + // Region 04 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK22-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK22-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK22-DAG: [[RVAR0]] = load i32*, i32** @d + // CK22-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 2 + // CK22-DAG: [[RVAR00]] = load i32*, i32** @d + + // CK22: call void [[CALL04:@.+]](i32* {{[^,]+}}) + #pragma omp target map(d[2:5]) + { d[3]+=1; } + + // Region 05 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[ST]]* @sa to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[ST]]* @sa to i8*), i8** [[P0]] + + // CK22: call void [[CALL05:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(sa) + { sa.fa+=1; } + + // Region 06 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[P0]] + + // CK22: call void [[CALL06:@.+]]([100 x [[ST]]]* {{[^,]+}}) + #pragma omp target map(sc) + { sc[3].fa+=1; } + + // Region 07 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[ST]]** @sd to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[ST]]** @sd to i8*), i8** [[P0]] + + // CK22: call void [[CALL07:@.+]]([[ST]]** {{[^,]+}}) + #pragma omp target map(sd) + { sd[3].fa+=1; } + + // Region 08 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[ST]]* getelementptr inbounds ([100 x [[ST]]], [100 x [[ST]]]* @sc, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]] + + // CK22: call void [[CALL08:@.+]]([100 x [[ST]]]* {{[^,]+}}) + #pragma omp target map(sc[1:4]) + { sc[3].fa+=1; } + + // Region 09 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK22-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[RVAR0:%.+]] to i8* + // CK22-DAG: [[CPVAL0]] = bitcast [[ST]]* [[SEC0:%.+]] to i8* + // CK22-DAG: [[RVAR0]] = load [[ST]]*, [[ST]]** @sd + // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[RVAR00:%.+]], i{{.+}} 2 + // CK22-DAG: [[RVAR00]] = load [[ST]]*, [[ST]]** @sd + + // CK22: call void [[CALL09:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(sd[2:5]) + { sd[3].fa+=1; } + + // Region 10 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[STT]]* @sta to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[STT]]* @sta to i8*), i8** [[P0]] + + // CK22: call void [[CALL10:@.+]]([[STT]]* {{[^,]+}}) + #pragma omp target map(sta) + { sta.fa+=1; } + + // Region 11 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[P0]] + + // CK22: call void [[CALL11:@.+]]([100 x [[STT]]]* {{[^,]+}}) + #pragma omp target map(stc) + { stc[3].fa+=1; } + + // Region 12 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[STT]]** @std to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[STT]]** @std to i8*), i8** [[P0]] + + // CK22: call void [[CALL12:@.+]]([[STT]]** {{[^,]+}}) + #pragma omp target map(std) + { std[3].fa+=1; } + + // Region 13 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[STT]]* getelementptr inbounds ([100 x [[STT]]], [100 x [[STT]]]* @stc, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]] + + // CK22: call void [[CALL13:@.+]]([100 x [[STT]]]* {{[^,]+}}) + #pragma omp target map(stc[1:4]) + { stc[3].fa+=1; } + + // Region 14 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK22-DAG: [[CBPVAL0]] = bitcast [[STT]]* [[RVAR0:%.+]] to i8* + // CK22-DAG: [[CPVAL0]] = bitcast [[STT]]* [[SEC0:%.+]] to i8* + // CK22-DAG: [[RVAR0]] = load [[STT]]*, [[STT]]** @std + // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[STT]]* [[RVAR00:%.+]], i{{.+}} 2 + // CK22-DAG: [[RVAR00]] = load [[STT]]*, [[STT]]** @std + + // CK22: call void [[CALL14:@.+]]([[STT]]* {{[^,]+}}) + #pragma omp target map(std[2:5]) + { std[3].fa+=1; } + + return 0; +} +// CK22: define {{.+}}[[CALL00]] +// CK22: define {{.+}}[[CALL01]] +// CK22: define {{.+}}[[CALL02]] +// CK22: define {{.+}}[[CALL03]] +// CK22: define {{.+}}[[CALL04]] +// CK22: define {{.+}}[[CALL05]] +// CK22: define {{.+}}[[CALL06]] +// CK22: define {{.+}}[[CALL07]] +// CK22: define {{.+}}[[CALL08]] +// CK22: define {{.+}}[[CALL09]] +// CK22: define {{.+}}[[CALL10]] +// CK22: define {{.+}}[[CALL11]] +// CK22: define {{.+}}[[CALL12]] +// CK22: define {{.+}}[[CALL13]] +// CK22: define {{.+}}[[CALL14]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -std=c++11 -DCK23 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64 +// RUN: %clang_cc1 -std=c++11 -DCK23 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -std=c++11 -fopenmp -fomptargets=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 CK23 --check-prefix CK23-64 +// RUN: %clang_cc1 -std=c++11 -DCK23 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32 +// RUN: %clang_cc1 -std=c++11 -DCK23 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -std=c++11 -fopenmp -fomptargets=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 CK23 --check-prefix CK23-32 +#ifdef CK23 + +// CK23: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23-LABEL: explicit_maps_inside_captured +int explicit_maps_inside_captured(int a){ + float b; + float c[100]; + float *d; + + // CK23: call void @{{.*}}explicit_maps_inside_captured{{.*}}([[SA:%.+]]* {{.*}}) + // CK23: define {{.*}}explicit_maps_inside_captured{{.*}} + [&](void){ + // Region 00 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast i32* [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load i32*, i32** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load i32*, i32** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(a) + { a+=1; } + // Region 01 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float* [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load float*, float** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load float*, float** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL01:@.+]](float* {{[^,]+}}) + #pragma omp target map(b) + { b+=1; } + // Region 02 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast [100 x float]* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast [100 x float]* [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load [100 x float]*, [100 x float]** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load [100 x float]*, [100 x float]** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL02:@.+]]([100 x float]* {{[^,]+}}) + #pragma omp target map(c) + { c[3]+=1; } + + // Region 03 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast float** [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float** [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL03:@.+]](float** {{[^,]+}}) + #pragma omp target map(d) + { d[3]+=1; } + // Region 04 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast [100 x float]* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8* + // CK23-DAG: [[SEC0]] = getelementptr {{.*}}[100 x float]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2 + // CK23-DAG: [[VAR0]] = load [100 x float]*, [100 x float]** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load [100 x float]*, [100 x float]** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL04:@.+]]([100 x float]* {{[^,]+}}) + #pragma omp target map(c[2:4]) + { c[3]+=1; } + + // Region 05 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast float* [[RVAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8* + // CK23-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]] + // CK23-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2 + // CK23-DAG: [[RVAR00]] = load float*, float** [[VAR00:%[^,]+]] + // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL05:@.+]](float* {{[^,]+}}) + #pragma omp target map(d[2:4]) + { d[3]+=1; } + }(); + return b; +} + +// CK23: define {{.+}}[[CALL00]] +// CK23: define {{.+}}[[CALL01]] +// CK23: define {{.+}}[[CALL02]] +// CK23: define {{.+}}[[CALL03]] +// CK23: define {{.+}}[[CALL04]] +// CK23: define {{.+}}[[CALL05]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-64 +// RUN: %clang_cc1 -DCK24 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK24 --check-prefix CK24-64 +// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-32 +// RUN: %clang_cc1 -DCK24 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK24 --check-prefix CK24-32 +#ifdef CK24 + +// CK24-DAG: [[SC:%.+]] = type { i32, [[SB:%.+]], [[SB:%.+]]*, [10 x i32] } +// CK24-DAG: [[SB]] = type { i32, [[SA:%.+]], [10 x [[SA:%.+]]], [10 x [[SA:%.+]]*], [[SA:%.+]]* } +// CK24-DAG: [[SA]] = type { i32, [[SA]]*, [10 x i32] } + +struct SA{ + int a; + struct SA *p; + int b[10]; +}; +struct SB{ + int a; + struct SA s; + struct SA sa[10]; + struct SA *sp[10]; + struct SA *p; +}; +struct SC{ + int a; + struct SB s; + struct SB *p; + int b[10]; +}; + +// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}] +// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}] +// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE07:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE08:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE09:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8] +// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE11:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}] +// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE12:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99] + +// CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}] +// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE16:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}] +// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE19:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE21:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE22:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE23:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}] +// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE24:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99] + +// CK24-LABEL: explicit_maps_struct_fields +int explicit_maps_struct_fields(int a){ + SC s; + SC *p; + +// Region 01 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + +// CK24: call void [[CALL01:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.a) + { s.a++; } + +// Region 02 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL02:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.s) + { s.a++; } + +// Region 03 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL03:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.s.a) + { s.a++; } + +// Region 04 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 3 + +// CK24: call void [[CALL04:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.b[:5]) + { s.a++; } + +// Region 05 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SB]]* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24: call void [[CALL05:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.p[:5]) + { s.a++; } + +// Region 06 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[10 x [[SA]]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL06:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.sa[3].a) + { s.a++; } + +// Region 07 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE07]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL07:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.sp[3]->a) + { s.a++; } + +// Region 08 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE08]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24: call void [[CALL08:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.p->a) + { s.a++; } + +// Region 09 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE09]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SB]]* [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL09:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.p->a) + { s.a++; } + +// Region 10 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SA]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL10:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.s.b[:2]) + { s.a++; } + +// Region 11 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE11]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[10 x i32]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC11]] = getelementptr {{.*}}[[SA]]* [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[SEC111]] = load [[SA]]*, [[SA]]** [[SEC1111:%[^,]+]], +// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + +// CK24: call void [[CALL11:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.s.p->b[:2]) + { s.a++; } + +// Region 12 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE12]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SA]]** [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] +// CK24-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] +// CK24-DAG: [[CBPVAL2]] = bitcast [[SA]]** [[SEC1]] to i8* +// CK24-DAG: [[CPVAL2]] = bitcast [[SA]]** [[SEC2:%.+]] to i8* +// CK24-DAG: [[SEC2]] = getelementptr {{.*}}[[SA]]* [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC22]] = load [[SA]]*, [[SA]]** [[SEC222:%[^,]+]], +// CK24-DAG: [[SEC222]] = getelementptr {{.*}}[[SB]]* [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC2222]] = load [[SB]]*, [[SB]]** [[SEC22222:%[^,]+]], +// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]] +// CK24-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]] +// CK24-DAG: [[CBPVAL3]] = bitcast [[SA]]** [[SEC2]] to i8* +// CK24-DAG: [[CPVAL3]] = bitcast i32* [[SEC3:%.+]] to i8* +// CK24-DAG: [[SEC3]] = getelementptr {{.*}}[[SA]]* [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC33]] = load [[SA]]*, [[SA]]** [[SEC333:%[^,]+]], +// CK24-DAG: [[SEC333]] = getelementptr {{.*}}[[SA]]* [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC3333]] = load [[SA]]*, [[SA]]** [[SEC33333:%[^,]+]], +// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}[[SB]]* [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC333333]] = load [[SB]]*, [[SB]]** [[SEC3333333:%[^,]+]], +// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + +// CK24: call void [[CALL12:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(s.p->p->p->a) + { s.a++; } + +// +// Same thing but starting from a pointer. +// +// Region 13 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 0 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL13:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->a) + { p->a++; } + +// Region 14 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL14:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.s) + { p->a++; } + +// Region 15 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL15:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.s.a) + { p->a++; } + +// Region 16 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE16]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE16]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 3 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL16:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->b[:5]) + { p->a++; } + +// Region 17 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE17]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE17]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SB]]* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL17:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->p[:5]) + { p->a++; } + +// Region 18 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE18]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE18]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[10 x [[SA]]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL18:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.sa[3].a) + { p->a++; } + +// Region 19 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE19]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE19]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL19:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.sp[3]->a) + { p->a++; } + +// Region 20 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE20]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE20]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL20:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->p->a) + { p->a++; } + +// Region 21 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE21]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE21]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SB]]* [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL21:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.p->a) + { p->a++; } + +// Region 22 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE22]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE22]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SA]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL22:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.s.b[:2]) + { p->a++; } + +// Region 23 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE23]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE23]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[10 x i32]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC11]] = getelementptr {{.*}}[[SA]]* [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[SEC111]] = load [[SA]]*, [[SA]]** [[SEC1111:%[^,]+]], +// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL23:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.p->b[:2]) + { p->a++; } + +// Region 24 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE24]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE24]]{{.+}}) +// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SA]]** [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 +// CK24-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] +// CK24-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] +// CK24-DAG: [[CBPVAL2]] = bitcast [[SA]]** [[SEC1]] to i8* +// CK24-DAG: [[CPVAL2]] = bitcast [[SA]]** [[SEC2:%.+]] to i8* +// CK24-DAG: [[SEC2]] = getelementptr {{.*}}[[SA]]* [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC22]] = load [[SA]]*, [[SA]]** [[SEC222:%[^,]+]], +// CK24-DAG: [[SEC222]] = getelementptr {{.*}}[[SB]]* [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC2222]] = load [[SB]]*, [[SB]]** [[SEC22222:%[^,]+]], +// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}[[SC]]* [[VAR0000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 +// CK24-DAG: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]] +// CK24-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]] +// CK24-DAG: [[CBPVAL3]] = bitcast [[SA]]** [[SEC2]] to i8* +// CK24-DAG: [[CPVAL3]] = bitcast i32* [[SEC3:%.+]] to i8* +// CK24-DAG: [[SEC3]] = getelementptr {{.*}}[[SA]]* [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC33]] = load [[SA]]*, [[SA]]** [[SEC333:%[^,]+]], +// CK24-DAG: [[SEC333]] = getelementptr {{.*}}[[SA]]* [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC3333]] = load [[SA]]*, [[SA]]** [[SEC33333:%[^,]+]], +// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}[[SB]]* [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC333333]] = load [[SB]]*, [[SB]]** [[SEC3333333:%[^,]+]], +// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}[[SC]]* [[VAR00000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR0000]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL24:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->p->p->p->a) + { p->a++; } + + return s.a; +} + +// CK24: define {{.+}}[[CALL01]] +// CK24: define {{.+}}[[CALL02]] +// CK24: define {{.+}}[[CALL03]] +// CK24: define {{.+}}[[CALL04]] +// CK24: define {{.+}}[[CALL05]] +// CK24: define {{.+}}[[CALL06]] +// CK24: define {{.+}}[[CALL07]] +// CK24: define {{.+}}[[CALL08]] +// CK24: define {{.+}}[[CALL09]] +// CK24: define {{.+}}[[CALL10]] +// CK24: define {{.+}}[[CALL11]] +// CK24: define {{.+}}[[CALL12]] +// CK24: define {{.+}}[[CALL13]] +// CK24: define {{.+}}[[CALL14]] +// CK24: define {{.+}}[[CALL15]] +// CK24: define {{.+}}[[CALL16]] +// CK24: define {{.+}}[[CALL17]] +// CK24: define {{.+}}[[CALL18]] +// CK24: define {{.+}}[[CALL19]] +// CK24: define {{.+}}[[CALL20]] +// CK24: define {{.+}}[[CALL21]] +// CK24: define {{.+}}[[CALL22]] +// CK24: define {{.+}}[[CALL23]] +// CK24: define {{.+}}[[CALL24]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK25 -std=c++11 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK25 --check-prefix CK25-64 +// RUN: %clang_cc1 -DCK25 -std=c++11 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -std=c++11 -fomptargets=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 CK25 --check-prefix CK25-64 +// RUN: %clang_cc1 -DCK25 -std=c++11 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK25 --check-prefix CK25-32 +// RUN: %clang_cc1 -DCK25 -std=c++11 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -std=c++11 -fomptargets=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 CK25 --check-prefix CK25-32 +#ifdef CK25 +// CK25: [[ST:%.+]] = type { i32, float } +// CK25: [[CA00:%.+]] = type { [[ST]]* } +// CK25: [[CA01:%.+]] = type { i32* } + +// CK25: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK25: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK25: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK25-LABEL: explicit_maps_with_inner_lambda + +template <int X, typename T> +struct CC { + T A; + float B; + + int foo(T arg) { + // Region 00 + // CK25-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK25-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK25-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK25-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK25-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK25-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK25-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK25-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK25-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK25-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0:%.+]], i{{.+}} 0, i{{.+}} 0 + + // CK25: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(to:A) + { + [&]() { + A += 1; + }(); + } + + // Region 01 + // CK25-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK25-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK25-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK25-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK25-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK25-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK25-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK25-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK25-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK25: call void [[CALL01:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:arg) + { + [&]() { + arg += 1; + }(); + } + + return A+arg; + } +}; + +int explicit_maps_with_inner_lambda(int a){ + CC<123,int> c; + return c.foo(a); +} + +// CK25: define {{.+}}[[CALL00]]([[ST]]* [[VAL:%.+]]) +// CK25: store [[ST]]* [[VAL]], [[ST]]** [[VALADDR:%[^,]+]], +// CK25: [[VAL1:%.+]] = load [[ST]]*, [[ST]]** [[VALADDR]], +// CK25: [[VALADDR1:%.+]] = getelementptr inbounds [[CA00]], [[CA00]]* [[CA:%[^,]+]], i32 0, i32 0 +// CK25: store [[ST]]* [[VAL1]], [[ST]]** [[VALADDR1]], +// CK25: call void {{.*}}[[LAMBDA:@.+]]{{.*}}([[CA00]]* [[CA]]) + +// CK25: define {{.+}}[[LAMBDA]] + +// CK25: define {{.+}}[[CALL01]](i32* {{.*}}[[VAL:%.+]]) +// CK25: store i32* [[VAL]], i32** [[VALADDR:%[^,]+]], +// CK25: [[VAL1:%.+]] = load i32*, i32** [[VALADDR]], +// CK25: [[VALADDR1:%.+]] = getelementptr inbounds [[CA01]], [[CA01]]* [[CA:%[^,]+]], i32 0, i32 0 +// CK25: store i32* [[VAL1]], i32** [[VALADDR1]], +// CK25: call void {{.*}}[[LAMBDA]]{{.*}}([[CA01]]* [[CA]]) +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK26 -std=c++11 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK26 --check-prefix CK26-64 +// RUN: %clang_cc1 -DCK26 -std=c++11 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -std=c++11 -fomptargets=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 CK26 --check-prefix CK26-64 +// RUN: %clang_cc1 -DCK26 -std=c++11 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK26 --check-prefix CK26-32 +// RUN: %clang_cc1 -DCK26 -std=c++11 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -std=c++11 -fomptargets=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 CK26 --check-prefix CK26-32 +#ifdef CK26 +// CK26: [[ST:%.+]] = type { i32, float*, i32, float* } + +// CK26: [[SIZE00:@.+]] = private {{.*}}constant [2 x i[[Z:64|32]]] [i[[Z:64|32]] {{32|16}}, i[[Z:64|32]] 4] +// CK26: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3] + +// CK26: [[SIZE01:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4] +// CK26: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3] + +// CK26: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4] +// CK26: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3] + +// CK26: [[SIZE03:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4] +// CK26: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3] + +// CK26-LABEL: explicit_maps_with_private_class_members + +struct CC { + int fA; + float &fB; + int pA; + float &pB; + + CC(float &B) : fB(B), pB(B) { + + // CK26: call {{.*}}@__kmpc_fork_call{{.*}} [[OUTCALL:@.+]] to void (i32*, i32*, ...)* + // define {{.*}}void [[OUTCALL]] + #pragma omp parallel firstprivate(fA,fB) private(pA,pB) + { + // Region 00 + // CK26-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK26-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK26-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK26-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK26-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK26-DAG: [[CPVAL0]] = bitcast [[ST]]* [[VAR0]] to i8* + + // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK26-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK26-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK26-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK26-DAG: [[VAR1]] = load i32*, i32** [[PVT:%.+]], + // CK26-DAG: [[SEC1]] = load i32*, i32** [[PVT]], + + // CK26: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(fA) + { + ++fA; + } + + // Region 01 + // CK26-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK26-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK26-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK26-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK26-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK26-DAG: [[CPVAL0]] = bitcast [[ST]]* [[VAR0]] to i8* + + // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK26-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK26-DAG: [[CBPVAL1]] = bitcast float* [[VAR1:%.+]] to i8* + // CK26-DAG: [[CPVAL1]] = bitcast float* [[SEC1:%.+]] to i8* + // CK26-DAG: [[VAR1]] = load float*, float** [[PVT:%.+]], + // CK26-DAG: [[SEC1]] = load float*, float** [[PVT]], + + // CK26: call void [[CALL01:@.+]]([[ST]]* {{[^,]+}}, float* {{[^,]+}}) + #pragma omp target map(fB) + { + fB += 1.0; + } + + // Region 02 + // CK26-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK26-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK26-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK26-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK26-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK26-DAG: [[CPVAL0]] = bitcast [[ST]]* [[VAR0]] to i8* + + // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK26-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK26-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK26-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK26-DAG: [[VAR1]] = load i32*, i32** [[PVT:%.+]], + // CK26-DAG: [[SEC1]] = load i32*, i32** [[PVT]], + + // CK26: call void [[CALL02:@.+]]([[ST]]* {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(pA) + { + ++pA; + } + + // Region 01 + // CK26-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK26-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK26-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK26-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK26-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK26-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK26-DAG: [[CPVAL0]] = bitcast [[ST]]* [[VAR0]] to i8* + + // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK26-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK26-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK26-DAG: [[CBPVAL1]] = bitcast float* [[VAR1:%.+]] to i8* + // CK26-DAG: [[CPVAL1]] = bitcast float* [[SEC1:%.+]] to i8* + // CK26-DAG: [[VAR1]] = load float*, float** [[PVT:%.+]], + // CK26-DAG: [[SEC1]] = load float*, float** [[PVT]], + + // CK26: call void [[CALL03:@.+]]([[ST]]* {{[^,]+}}, float* {{[^,]+}}) + #pragma omp target map(pB) + { + pB += 1.0; + } + } + } + + int foo() { + return fA + pA; + } +}; + +// Make sure the private instance is used in all target regions. +// CK26: define {{.+}}[[CALL00]]({{.*}}i32*{{.*}}[[PVTARG:%.+]]) +// CK26: store i32* [[PVTARG]], i32** [[PVTADDR:%.+]], +// CK26: [[ADDR:%.+]] = load i32*, i32** [[PVTADDR]], +// CK26: [[VAL:%.+]] = load i32, i32* [[ADDR]], +// CK26: add nsw i32 [[VAL]], 1 + +// CK26: define {{.+}}[[CALL01]]({{.*}}float*{{.*}}[[PVTARG:%.+]]) +// CK26: store float* [[PVTARG]], float** [[PVTADDR:%.+]], +// CK26: [[ADDR:%.+]] = load float*, float** [[PVTADDR]], +// CK26: [[VAL:%.+]] = load float, float* [[ADDR]], +// CK26: [[EXT:%.+]] = fpext float [[VAL]] to double +// CK26: fadd double [[EXT]], 1.000000e+00 + +// CK26: define {{.+}}[[CALL02]]({{.*}}i32*{{.*}}[[PVTARG:%.+]]) +// CK26: store i32* [[PVTARG]], i32** [[PVTADDR:%.+]], +// CK26: [[ADDR:%.+]] = load i32*, i32** [[PVTADDR]], +// CK26: [[VAL:%.+]] = load i32, i32* [[ADDR]], +// CK26: add nsw i32 [[VAL]], 1 + +// CK26: define {{.+}}[[CALL03]]({{.*}}float*{{.*}}[[PVTARG:%.+]]) +// CK26: store float* [[PVTARG]], float** [[PVTADDR:%.+]], +// CK26: [[ADDR:%.+]] = load float*, float** [[PVTADDR]], +// CK26: [[VAL:%.+]] = load float, float* [[ADDR]], +// CK26: [[EXT:%.+]] = fpext float [[VAL]] to double +// CK26: fadd double [[EXT]], 1.000000e+00 + +int explicit_maps_with_private_class_members(){ + float B; + CC c(B); + return c.foo(); +} +#endif #endif |