clang/test/OpenMP/target_map_codegen_21.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64
   7 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64
   9 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32
  10 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32
  12
  13 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64
  14 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64
  16 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32
  17 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32
  19
  20 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64
  21 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64
  23 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32
  24 // RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32
  26
  27 // RUN: %clang_cc1 -DCK22 -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-ONLY21 %s
  28 // RUN: %clang_cc1 -DCK22 -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-ONLY21 %s
  30 // RUN: %clang_cc1 -DCK22 -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-ONLY21 %s
  31 // RUN: %clang_cc1 -DCK22 -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-ONLY21 %s
  33 // SIMD-ONLY21-NOT: {{__kmpc|__tgt}}
  34 #ifdef CK22
  35
  36 // CK22-DAG: [[ST:%.+]] = type { float }
  37 // CK22-DAG: [[STT:%.+]] = type { i32 }
  38
  39 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  40 // CK22: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
  41 // CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  42
  43 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  44 // CK22: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
  45 // CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  46
  47 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  48 // CK22: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
  49 // CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  50
  51 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  52 // CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
  53 // CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  54
  55 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  56 // CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
  57 // CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
  58
  59 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  60 // CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
  61 // CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  62
  63 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  64 // CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
  65 // CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  66
  67 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  68 // CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
  69 // CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  70
  71 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  72 // CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
  73 // CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  74
  75 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  76 // CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
  77 // CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
  78
  79 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  80 // CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
  81 // CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  82
  83 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  84 // CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
  85 // CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  86
  87 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  88 // CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
  89 // CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  90
  91 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  92 // CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
  93 // CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  94
  95 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  96 // CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
  97 // CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
  98
  99 int a;
 100 int c[100];
 101 int *d;
 102
 103 struct ST {
 104   float fa;
 105 };
 106
 107 ST sa ;
 108 ST sc[100];
 109 ST *sd;
 110
 111 template<typename T>
 112 struct STT {
 113   T fa;
 114 };
 115
 116 STT<int> sta ;
 117 STT<int> stc[100];
 118 STT<int> *std;
 119
 120 // CK22-LABEL: explicit_maps_globals{{.*}}(
 121 int explicit_maps_globals(void){
 122 // Region 00
 123 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 124 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 125 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 126 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 127 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 128 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 129 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 130
 131 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 132 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 133 // CK22-DAG: store ptr @a, ptr [[BP0]]
 134 // CK22-DAG: store ptr @a, ptr [[P0]]
 135
 136 // CK22: call void [[CALL00:@.+]](ptr {{[^,]+}})
 137 #pragma omp target map(a)
 138   { a+=1; }
 139
 140 // Region 01
 141 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 142 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 143 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 144 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 145 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 146 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 147 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 148
 149 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 150 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 151 // CK22-DAG: store ptr @c, ptr [[BP0]]
 152 // CK22-DAG: store ptr @c, ptr [[P0]]
 153
 154 // CK22: call void [[CALL01:@.+]](ptr {{[^,]+}})
 155 #pragma omp target map(c)
 156   { c[3]+=1; }
 157
 158 // Region 02
 159 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 160 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 161 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 162 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 163 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 164 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 165 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 166
 167 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 168 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 169 // CK22-DAG: store ptr @d, ptr [[BP0]]
 170 // CK22-DAG: store ptr @d, ptr [[P0]]
 171
 172 // CK22: call void [[CALL02:@.+]](ptr {{[^,]+}})
 173 #pragma omp target map(d)
 174   { d[3]+=1; }
 175
 176 // Region 03
 177 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 178 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 179 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 180 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 181 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 182 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 183 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 184
 185 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 186 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 187 // CK22-DAG: store ptr @c, ptr [[BP0]]
 188 // CK22-DAG: store ptr getelementptr inbounds ([100 x i32], ptr @c, i{{.+}} 0, i{{.+}} 1), ptr [[P0]]
 189
 190 // CK22: call void [[CALL03:@.+]](ptr {{[^,]+}})
 191 #pragma omp target map(c [1:4])
 192   { c[3]+=1; }
 193
 194 // Region 04
 195 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 196 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 197 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 198 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 199 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 200 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 201 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 202
 203 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 204 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 205 // CK22-DAG: store ptr @d, ptr [[BP0]]
 206 // CK22-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 207 // CK22-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
 208 // CK22-DAG: [[RVAR00]] = load ptr, ptr @d
 209
 210 // CK22: call void [[CALL04:@.+]](ptr {{[^,]+}})
 211 #pragma omp target map(d [2:5])
 212   { d[3]+=1; }
 213
 214 // Region 05
 215 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 216 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 217 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 218 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 219 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 220 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 221 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 222
 223 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 224 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 225 // CK22-DAG: store ptr @sa, ptr [[BP0]]
 226 // CK22-DAG: store ptr @sa, ptr [[P0]]
 227
 228 // CK22: call void [[CALL05:@.+]](ptr {{[^,]+}})
 229 #pragma omp target map(sa)
 230   { sa.fa+=1; }
 231
 232 // Region 06
 233 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 234 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 235 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 236 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 237 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 238 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 239 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 240
 241 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 242 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 243 // CK22-DAG: store ptr @sc, ptr [[BP0]]
 244 // CK22-DAG: store ptr @sc, ptr [[P0]]
 245
 246 // CK22: call void [[CALL06:@.+]](ptr {{[^,]+}})
 247 #pragma omp target map(sc)
 248   { sc[3].fa+=1; }
 249
 250 // Region 07
 251 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 252 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 253 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 254 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 255 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 256 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 257 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 258
 259 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 260 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 261 // CK22-DAG: store ptr @sd, ptr [[BP0]]
 262 // CK22-DAG: store ptr @sd, ptr [[P0]]
 263
 264 // CK22: call void [[CALL07:@.+]](ptr {{[^,]+}})
 265 #pragma omp target map(sd)
 266   { sd[3].fa+=1; }
 267
 268 // Region 08
 269 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 270 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 271 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 272 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 273 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 274 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 275 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 276
 277 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 278 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 279 // CK22-DAG: store ptr @sc, ptr [[BP0]]
 280 // CK22-DAG: store ptr getelementptr inbounds ([100 x [[ST]]], ptr @sc, i{{.+}} 0, i{{.+}} 1), ptr [[P0]]
 281
 282 // CK22: call void [[CALL08:@.+]](ptr {{[^,]+}})
 283 #pragma omp target map(sc [1:4])
 284   { sc[3].fa+=1; }
 285
 286 // Region 09
 287 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 288 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 289 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 290 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 291 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 292 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 293 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 294
 295 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 296 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 297 // CK22-DAG: store ptr @sd, ptr [[BP0]]
 298 // CK22-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 299 // CK22-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
 300 // CK22-DAG: [[RVAR00]] = load ptr, ptr @sd
 301
 302 // CK22: call void [[CALL09:@.+]](ptr {{[^,]+}})
 303 #pragma omp target map(sd [2:5])
 304   { sd[3].fa+=1; }
 305
 306 // Region 10
 307 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 308 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 309 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 310 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 311 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 312 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 313 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 314
 315 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 316 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 317 // CK22-DAG: store ptr @sta, ptr [[BP0]]
 318 // CK22-DAG: store ptr @sta, ptr [[P0]]
 319
 320 // CK22: call void [[CALL10:@.+]](ptr {{[^,]+}})
 321 #pragma omp target map(sta)
 322   { sta.fa+=1; }
 323
 324 // Region 11
 325 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 326 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 327 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 328 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 329 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 330 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 331 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 332
 333 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 334 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 335 // CK22-DAG: store ptr @stc, ptr [[BP0]]
 336 // CK22-DAG: store ptr @stc, ptr [[P0]]
 337
 338 // CK22: call void [[CALL11:@.+]](ptr {{[^,]+}})
 339 #pragma omp target map(stc)
 340   { stc[3].fa+=1; }
 341
 342 // Region 12
 343 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 344 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 345 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 346 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 347 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 348 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 349 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 350
 351 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 352 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 353 // CK22-DAG: store ptr @std, ptr [[BP0]]
 354 // CK22-DAG: store ptr @std, ptr [[P0]]
 355
 356 // CK22: call void [[CALL12:@.+]](ptr {{[^,]+}})
 357 #pragma omp target map(std)
 358   { std[3].fa+=1; }
 359
 360 // Region 13
 361 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 362 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 363 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 364 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 365 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 366 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 367 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 368
 369 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 370 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 371 // CK22-DAG: store ptr @stc, ptr [[BP0]]
 372 // CK22-DAG: store ptr getelementptr inbounds ([100 x [[STT]]], ptr @stc, i{{.+}} 0, i{{.+}} 1),  ptr [[P0]]
 373
 374 // CK22: call void [[CALL13:@.+]](ptr {{[^,]+}})
 375 #pragma omp target map(stc [1:4])
 376   { stc[3].fa+=1; }
 377
 378 // Region 14
 379 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 380 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 381 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 382 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 383 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 384 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 385 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 386
 387 // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 388 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 389 // CK22-DAG: store ptr @std, ptr [[BP0]]
 390 // CK22-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 391 // CK22-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
 392 // CK22-DAG: [[RVAR00]] = load ptr, ptr @std
 393
 394 // CK22: call void [[CALL14:@.+]](ptr {{[^,]+}})
 395 #pragma omp target map(std [2:5])
 396   { std[3].fa+=1; }
 397
 398   return 0;
 399 }
 400 // CK22: define {{.+}}[[CALL00]]
 401 // CK22: define {{.+}}[[CALL01]]
 402 // CK22: define {{.+}}[[CALL02]]
 403 // CK22: define {{.+}}[[CALL03]]
 404 // CK22: define {{.+}}[[CALL04]]
 405 // CK22: define {{.+}}[[CALL05]]
 406 // CK22: define {{.+}}[[CALL06]]
 407 // CK22: define {{.+}}[[CALL07]]
 408 // CK22: define {{.+}}[[CALL08]]
 409 // CK22: define {{.+}}[[CALL09]]
 410 // CK22: define {{.+}}[[CALL10]]
 411 // CK22: define {{.+}}[[CALL11]]
 412 // CK22: define {{.+}}[[CALL12]]
 413 // CK22: define {{.+}}[[CALL13]]
 414 // CK22: define {{.+}}[[CALL14]]
 415 #endif // CK22
 416 #endif