clang/test/OpenMP/target_map_codegen_23.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   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-64
   7 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
   8 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK24 --check-prefix CK24-64
   9 // 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-32
  10 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  11 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK24 --check-prefix CK24-32
  12
  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-64
  14 // 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 %s
  15 // 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-64
  16 // 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-32
  17 // 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 %s
  18 // 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-32
  19
  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-64
  21 // 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 %s
  22 // 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-64
  23 // 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-32
  24 // 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 %s
  25 // 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-32
  26
  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 %s
  28 // 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 %s
  29 // 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 %s
  30 // 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 %s
  31 // 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 %s
  32 // 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 %s
  33 // SIMD-ONLY23-NOT: {{__kmpc|__tgt}}
  34 #ifdef CK24
  35
  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
  40 struct SA{
  41   int a;
  42   struct SA *p;
  43   int b[10];
  44 };
  45 struct SB{
  46   int a;
  47   struct SA s;
  48   struct SA sa[10];
  49   struct SA *sp[10];
  50   struct SA *p;
  51 };
  52 struct 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 0
  60 // 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 0
  64 // 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 0
  68 // 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 0
  72 // 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 0
  76 // 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 0
  80 // 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 0
  84 // 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 0
  88 // 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 0
  92 // 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 0
  96 // 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 0
 100 // 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 0
 104 // 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 0
 108 // 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_fields
 112 int explicit_maps_struct_fields(int a){
 113   SC s;
 114   SC *p;
 115
 116 // Region 01
 117 // 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 2
 119 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 120 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 121 // 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{{.+}} 0
 126 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 127 // 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{{.+}} 0
 130
 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 13
 139 // 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 2
 141 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 142 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 143 // 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{{.+}} 0
 148 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 149 // 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{{.+}} 0
 152
 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 14
 161 // 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 2
 163 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 164 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 165 // 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{{.+}} 0
 170 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 171 // 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{{.+}} 1
 174 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 175
 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 15
 184 // 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 2
 186 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 187 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 188 // 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{{.+}} 0
 193 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 194 // 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{{.+}} 0
 197 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
 198 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 199
 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 16
 208 // 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 2
 210 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 211 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 212 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 213 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 214 // 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{{.+}} 0
 220 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 221 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 222 // 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{{.+}} 0
 226 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 3
 227
 228 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 229 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 230 // 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 17
 241 // 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 2
 243 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 244 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 245 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 246 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 247 // 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{{.+}} 0
 253 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 254 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 255 // 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{{.+}} 2
 259
 260 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 261 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 262 // CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]
 263 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 264 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
 265 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 266 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
 267
 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 18
 277 // 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 2
 279 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 280 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 281 // 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{{.+}} 0
 286 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 287 // 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{{.+}} 0
 290 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
 291 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
 292 // CK24-DAG: [[SEC0000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 293
 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 19
 302 // 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 2
 304 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 305 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 306 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 307 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 308 // 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{{.+}} 0
 314 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 315 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 316 // 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{{.+}} 3
 320 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
 321 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 322
 323 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 324 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 325 // 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{{.+}} 2
 329 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
 330 // 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{{.+}} 0
 333 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 334 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
 335 // CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
 336 // CK24-DAG: [[SEC11111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
 337
 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 20
 347 // 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 2
 349 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 350 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 351 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 352 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 353 // 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{{.+}} 0
 359 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 360 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 361 // 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{{.+}} 2
 365
 366 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 367 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 368 // CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]
 369 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 370 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
 371 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 372 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
 373
 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 21
 383 // 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 2
 385 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 386 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 387 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 388 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 389 // 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{{.+}} 0
 395 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 396 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 397 // 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{{.+}} 4
 401 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 402
 403 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 404 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 405 // 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{{.+}} 2
 409 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
 410 // CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]
 411 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]
 412 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
 413 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 414 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4
 415 // CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
 416
 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 22
 426 // 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 2
 428 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 429 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 430 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 431 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 432 // 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{{.+}} 0
 438 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 439 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 440 // 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{{.+}} 0
 444 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
 445 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
 446 // CK24-DAG: [[SEC0000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 447
 448 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 449 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 450 // 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 23
 461 // 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 2
 463 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 464 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 465 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 466 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 467 // 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{{.+}} 0
 473 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 474 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 475 // 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{{.+}} 4
 479 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
 480
 481 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 482 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 483 // 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{{.+}} 2
 487 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
 488 // 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{{.+}} 0
 491 // CK24-DAG: [[SEC11]] = getelementptr {{.*}}ptr [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2
 492 // CK24-DAG: [[SEC111]] = load ptr, ptr [[SEC1111:%[^,]+]],
 493 // CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4
 494 // CK24-DAG: [[SEC11111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
 495
 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 24
 505 // 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 2
 507 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 508 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 509 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 510 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 511 // 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{{.+}} 0
 517 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 518 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 519 // 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{{.+}} 2
 523
 524 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 525 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 526 // 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{{.+}} 4
 529 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
 530 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
 531
 532 // CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
 533 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
 534 // 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{{.+}} 1
 537 // CK24-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],
 538 // CK24-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4
 539 // CK24-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],
 540 // CK24-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[VAR0000:%.+]], i{{.+}} 0, i{{.+}} 2
 541
 542 // CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
 543 // CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
 544 // 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{{.+}} 0
 547 // CK24-DAG: [[SEC33]] = load ptr, ptr [[SEC333:%[^,]+]],
 548 // CK24-DAG: [[SEC333]] = getelementptr {{.*}}ptr [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1
 549 // CK24-DAG: [[SEC3333]] = load ptr, ptr [[SEC33333:%[^,]+]],
 550 // CK24-DAG: [[SEC33333]] = getelementptr {{.*}}ptr [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4
 551 // CK24-DAG: [[SEC333333]] = load ptr, ptr [[SEC3333333:%[^,]+]],
 552 // CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}ptr [[VAR00000:%.+]], i{{.+}} 0, i{{.+}} 2
 553
 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 // CK24
 581 #endif