582 lines · cpp
1// expected-no-diagnostics2#ifndef HEADER3#define HEADER4 5///==========================================================================///6// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-647// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s8// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-649// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-3210// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s11// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-3212 13// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-6414// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s15// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-6416// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-3217// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s18// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-3219 20// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-6421// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s22// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-6423// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-3224// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s25// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-3226 27// RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s28// RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s29// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s30// RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s31// RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s32// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s33// SIMD-ONLY23-NOT: {{__kmpc|__tgt}}34#ifdef CK2435 36// CK24-DAG: [[SC:%.+]] = type { i32, [[SB:%.+]], ptr, [10 x i32] }37// CK24-DAG: [[SB]] = type { i32, [[SA:%.+]], [10 x [[SA:%.+]]], [10 x ptr], ptr }38// CK24-DAG: [[SA]] = type { i32, ptr, [10 x i32] }39 40struct SA{41 int a;42 struct SA *p;43 int b[10];44};45struct SB{46 int a;47 struct SA s;48 struct SA sa[10];49 struct SA *sp[10];50 struct SA *p;51};52struct SC{53 int a;54 struct SB s;55 struct SB *p;56 int b[10];57};58 59// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 060// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 4]61// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]62 63// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 064// CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i64] [i64 4]65// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 35]66 67// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 068// CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 {{48|56}}]69// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 35]70 71// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 072// CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4]73// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 35]74 75// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 076// CK24: [[SIZE16:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 20]77// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659]78 79// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 080// CK24: [[SIZE17:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 {{3560|2880}}]81// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]82 83// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 084// CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i64] [i64 4]85// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i64] [i64 35]86 87// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 088// CK24: [[SIZE19:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 {{8|4}}, i64 4]89// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675]90 91// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 092// CK24: [[SIZE20:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 4]93// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]94 95// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 096// CK24: [[SIZE21:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 {{8|4}}, i64 4]97// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675]98 99// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0100// CK24: [[SIZE22:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 8]101// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659]102 103// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0104// CK24: [[SIZE23:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 {{8|4}}, i64 8]105// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675]106 107// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0108// CK24: [[SIZE24:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 {{8|4}}, i64 {{8|4}}, i64 4]109// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i64] [i64 32, i64 281474976710672, i64 16, i64 19]110 111// CK24-LABEL: explicit_maps_struct_fields112int explicit_maps_struct_fields(int a){113 SC s;114 SC *p;115 116// Region 01117// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])118// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2119// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]120// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3121// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]122// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]123// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]124 125// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0126// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0127// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]128// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]129// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0130 131// CK24: call void [[CALL01:@.+]](ptr {{[^,]+}})132#pragma omp target map(s.a)133 { s.a++; }134 135//136// Same thing but starting from a pointer.137//138// Region 13139// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])140// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2141// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]142// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3143// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]144// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]145// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]146 147// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0148// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0149// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]150// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]151// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 0152 153// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}154// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}155 156// CK24: call void [[CALL13:@.+]](ptr {{[^,]+}})157#pragma omp target map(p->a)158 { p->a++; }159 160// Region 14161// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])162// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2163// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]164// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3165// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]166// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]167// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]168 169// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0170// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0171// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]172// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]173// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1174// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1175 176// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}177// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}178 179// CK24: call void [[CALL14:@.+]](ptr {{[^,]+}})180#pragma omp target map(p->s.s)181 { p->a++; }182 183// Region 15184// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])185// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2186// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]187// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3188// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]189// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]190// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]191 192// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0193// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0194// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]195// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]196// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0197// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1198// CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1199 200// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}201// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}202 203// CK24: call void [[CALL15:@.+]](ptr {{[^,]+}})204#pragma omp target map(p->s.s.a)205 { p->a++; }206 207// Region 16208// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])209// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2210// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]211// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3212// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]213// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4214// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]215// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]216// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]217// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]218 219// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0220// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0221// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0222// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]223// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]224// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]225// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0226// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 3227 228// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1229// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1230// CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]231// CK24-DAG: store ptr [[SEC0]], ptr [[P1]]232 233// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}234// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}235 236// CK24: call void [[CALL16:@.+]](ptr {{[^,]+}})237#pragma omp target map(p->b[:5])238 { p->a++; }239 240// Region 17241// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])242// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2243// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]244// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3245// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]246// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4247// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]248// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]249// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]250// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]251 252// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0253// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0254// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0255// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]256// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]257// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]258// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2259 260// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1261// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1262// CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]263// CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]264// CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0265// CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],266// CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2267 268// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}269// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}270// CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}271 272// CK24: call void [[CALL17:@.+]](ptr {{[^,]+}})273#pragma omp target map(p->p[:5])274 { p->a++; }275 276// Region 18277// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])278// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2279// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]280// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3281// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]282// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]283// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]284 285// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0286// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0287// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]288// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]289// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0290// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3291// CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2292// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1293 294// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}295// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}296 297// CK24: call void [[CALL18:@.+]](ptr {{[^,]+}})298#pragma omp target map(p->s.sa[3].a)299 { p->a++; }300 301// Region 19302// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])303// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2304// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]305// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3306// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]307// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4308// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]309// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]310// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]311// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]312 313// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0314// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0315// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0316// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]317// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]318// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]319// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3320// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3321// CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1322 323// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1324// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1325// CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]326// CK24-DAG: store ptr [[SEC0]], ptr [[P1]]327 328// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2329// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2330// CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]331// CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]332// CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0333// CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],334// CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3335// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3336// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1337 338// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}339// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}340// CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}341 342// CK24: call void [[CALL19:@.+]](ptr {{[^,]+}})343#pragma omp target map(p->s.sp[3]->a)344 { p->a++; }345 346// Region 20347// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])348// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2349// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]350// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3351// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]352// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4353// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]354// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]355// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]356// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]357 358// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0359// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0360// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0361// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]362// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]363// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]364// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2365 366// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1367// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1368// CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]369// CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]370// CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0371// CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],372// CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2373 374// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}375// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}376// CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}377 378// CK24: call void [[CALL20:@.+]](ptr {{[^,]+}})379#pragma omp target map(p->p->a)380 { p->a++; }381 382// Region 21383// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])384// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2385// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]386// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3387// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]388// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4389// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]390// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]391// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]392// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]393 394// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0395// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0396// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0397// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]398// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]399// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]400// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4401// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1402 403// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1404// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1405// CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]406// CK24-DAG: store ptr [[SEC0]], ptr [[P1]]407 408// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2409// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2410// CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]411// CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]412// CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0413// CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],414// CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4415// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1416 417// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}418// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}419// CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}420 421// CK24: call void [[CALL21:@.+]](ptr {{[^,]+}})422#pragma omp target map(p->s.p->a)423 { p->a++; }424 425// Region 22426// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])427// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2428// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]429// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3430// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]431// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4432// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]433// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]434// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]435// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]436 437// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0438// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0439// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0440// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]441// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]442// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]443// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0444// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2445// CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1446// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1447 448// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1449// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1450// CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]451// CK24-DAG: store ptr [[SEC0]], ptr [[P1]]452 453// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}454// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}455 456// CK24: call void [[CALL22:@.+]](ptr {{[^,]+}})457#pragma omp target map(p->s.s.b[:2])458 { p->a++; }459 460// Region 23461// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])462// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2463// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]464// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3465// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]466// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4467// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]468// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]469// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]470// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]471 472// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0473// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0474// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0475// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]476// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]477// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]478// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4479// CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1480 481// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1482// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1483// CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]484// CK24-DAG: store ptr [[SEC0]], ptr [[P1]]485 486// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2487// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2488// CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]489// CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]490// CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0491// CK24-DAG: [[SEC11]] = getelementptr {{.*}}ptr [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2492// CK24-DAG: [[SEC111]] = load ptr, ptr [[SEC1111:%[^,]+]],493// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4494// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1495 496// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}497// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}498// CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}499 500// CK24: call void [[CALL23:@.+]](ptr {{[^,]+}})501#pragma omp target map(p->s.p->b[:2])502 { p->a++; }503 504// Region 24505// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])506// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2507// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]508// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3509// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]510// CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4511// CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]512// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]513// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]514// CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]515 516// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0517// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0518// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0519// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]520// CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]521// CK24-DAG: store i64 {{%.+}}, ptr [[S0]]522// CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2523 524// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1525// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1526// CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]527// CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]528// CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4529// CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],530// CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2531 532// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2533// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2534// CK24-DAG: store ptr [[SEC1]], ptr [[BP2]]535// CK24-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]536// CK24-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1537// CK24-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],538// CK24-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4539// CK24-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],540// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[VAR0000:%.+]], i{{.+}} 0, i{{.+}} 2541 542// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3543// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3544// CK24-DAG: store ptr [[SEC2]], ptr [[BP3]]545// CK24-DAG: store ptr [[SEC3:%.+]], ptr [[P3]]546// CK24-DAG: [[SEC3]] = getelementptr {{.*}}ptr [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0547// CK24-DAG: [[SEC33]] = load ptr, ptr [[SEC333:%[^,]+]],548// CK24-DAG: [[SEC333]] = getelementptr {{.*}}ptr [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1549// CK24-DAG: [[SEC3333]] = load ptr, ptr [[SEC33333:%[^,]+]],550// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}ptr [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4551// CK24-DAG: [[SEC333333]] = load ptr, ptr [[SEC3333333:%[^,]+]],552// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}ptr [[VAR00000:%.+]], i{{.+}} 0, i{{.+}} 2553 554// CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}555// CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}556// CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}557// CK24-DAG: [[VAR0000]] = load ptr, ptr %{{.+}}558// CK24-DAG: [[VAR00000]] = load ptr, ptr %{{.+}}559 560// CK24: call void [[CALL24:@.+]](ptr {{[^,]+}})561#pragma omp target map(p->p->p->p->a)562 { p->a++; }563 564 return s.a;565}566 567// CK24: define {{.+}}[[CALL01]]568// CK24: define {{.+}}[[CALL13]]569// CK24: define {{.+}}[[CALL14]]570// CK24: define {{.+}}[[CALL15]]571// CK24: define {{.+}}[[CALL16]]572// CK24: define {{.+}}[[CALL17]]573// CK24: define {{.+}}[[CALL18]]574// CK24: define {{.+}}[[CALL19]]575// CK24: define {{.+}}[[CALL20]]576// CK24: define {{.+}}[[CALL21]]577// CK24: define {{.+}}[[CALL22]]578// CK24: define {{.+}}[[CALL23]]579// CK24: define {{.+}}[[CALL24]]580#endif // CK24581#endif582