clang/test/OpenMP/target_map_codegen_21.cpp

// expected-no-diagnostics
#ifndef HEADER
#define HEADER

///==========================================================================///
// 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
// 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
// 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
// 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
// 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
// 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

// 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
// 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
// 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
// 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
// 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
// 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

// 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
// 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
// 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
// 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
// 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
// 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

// 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
// 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
// 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
// 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
// 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
// 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
// SIMD-ONLY21-NOT: {{__kmpc|__tgt}}
#ifdef CK22

// CK22-DAG: [[ST:%.+]] = type { float }
// CK22-DAG: [[STT:%.+]] = type { i32 }

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE00:@.+]] = private {{.*}}constant [2 x i64] [i64 4, i64 0]
// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE01:@.+]] = private {{.*}}constant [2 x i64] [i64 400, i64 0]
// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE02:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0]
// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE03:@.+]] = private {{.*}}constant [2 x i64] [i64 16, i64 0]
// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE04:@.+]] = private {{.*}}constant [3 x i64] [i64 20, i64 {{4|8}}, i64 0]
// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16384, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE05:@.+]] = private {{.*}}constant [2 x i64] [i64 4, i64 0]
// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE06:@.+]] = private {{.*}}constant [2 x i64] [i64 400, i64 0]
// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE07:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0]
// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE08:@.+]] = private {{.*}}constant [2 x i64] [i64 16, i64 0]
// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE09:@.+]] = private {{.*}}constant [3 x i64] [i64 20, i64 {{4|8}}, i64 0]
// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16384, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE10:@.+]] = private {{.*}}constant [2 x i64] [i64 4, i64 0]
// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE11:@.+]] = private {{.*}}constant [2 x i64] [i64 400, i64 0]
// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE12:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0]
// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE13:@.+]] = private {{.*}}constant [2 x i64] [i64 16, i64 0]
// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [2 x i64] [i64 35, i64 288]

// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE14:@.+]] = private {{.*}}constant [3 x i64] [i64 20, i64 {{4|8}}, i64 0]
// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16384, i64 288]

int a;
int c[100];
int *d;

struct ST {
  float fa;
};

ST sa ;
ST sc[100];
ST *sd;

template<typename T>
struct STT {
  T fa;
};

STT<int> sta ;
STT<int> stc[100];
STT<int> *std;

// CK22-LABEL: explicit_maps_globals{{.*}}(
int explicit_maps_globals(void){
// Region 00
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @a, ptr [[BP0]]
// CK22-DAG: store ptr @a, ptr [[P0]]

// CK22: call void [[CALL00:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(a)
  { a+=1; }

// Region 01
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @c, ptr [[BP0]]
// CK22-DAG: store ptr @c, ptr [[P0]]

// CK22: call void [[CALL01:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(c)
  { c[3]+=1; }

// Region 02
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @d, ptr [[BP0]]
// CK22-DAG: store ptr @d, ptr [[P0]]

// CK22: call void [[CALL02:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(d)
  { d[3]+=1; }

// Region 03
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @c, ptr [[BP0]]
// CK22-DAG: store ptr getelementptr inbounds nuw (i8, ptr @c, i{{32|64}} 4), ptr [[P0]]

// CK22: call void [[CALL03:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(c [1:4])
  { c[3]+=1; }

// Region 04

//  &d[0], &d[2], 5 * sizeof(d[0]), TO | FROM | PARAM
//  &d,    &d[2], sizeof(d),        ATTACH

// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
// CK22-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
// CK22-DAG: [[RVAR0]] = load ptr, ptr @d
// CK22-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR00]] = load ptr, ptr @d

// CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
// CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CK22-DAG: store ptr @d, ptr [[BP1]]
// CK22-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
// CK22-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[RVAR1:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR1]] = load ptr, ptr @d

// CK22: call void [[CALL04:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(d [2:5])
  { d[3]+=1; }

// Region 05
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @sa, ptr [[BP0]]
// CK22-DAG: store ptr @sa, ptr [[P0]]

// CK22: call void [[CALL05:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(sa)
  { sa.fa+=1; }

// Region 06
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @sc, ptr [[BP0]]
// CK22-DAG: store ptr @sc, ptr [[P0]]

// CK22: call void [[CALL06:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(sc)
  { sc[3].fa+=1; }

// Region 07
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @sd, ptr [[BP0]]
// CK22-DAG: store ptr @sd, ptr [[P0]]

// CK22: call void [[CALL07:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(sd)
  { sd[3].fa+=1; }

// Region 08
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @sc, ptr [[BP0]]
// CK22-DAG: store ptr getelementptr inbounds nuw (i8, ptr @sc, i{{32|64}} 4), ptr [[P0]]

// CK22: call void [[CALL08:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(sc [1:4])
  { sc[3].fa+=1; }

// Region 09

//  &sd[0], &sd[2], 5 * sizeof(sd[0]), TO | FROM | ATTACH
//  &sd,    &sd[2], sizeof(sd),        ATTACH

// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
// CK22-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
// CK22-DAG: [[RVAR0]] = load ptr, ptr @sd
// CK22-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR00]] = load ptr, ptr @sd

// CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
// CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CK22-DAG: store ptr @sd, ptr [[BP1]]
// CK22-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
// CK22-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[RVAR1:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR1]] = load ptr, ptr @sd

// CK22: call void [[CALL09:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(sd [2:5])
  { sd[3].fa+=1; }

// Region 10
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @sta, ptr [[BP0]]
// CK22-DAG: store ptr @sta, ptr [[P0]]

// CK22: call void [[CALL10:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(sta)
  { sta.fa+=1; }

// Region 11
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @stc, ptr [[BP0]]
// CK22-DAG: store ptr @stc, ptr [[P0]]

// CK22: call void [[CALL11:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(stc)
  { stc[3].fa+=1; }

// Region 12
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @std, ptr [[BP0]]
// CK22-DAG: store ptr @std, ptr [[P0]]

// CK22: call void [[CALL12:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(std)
  { std[3].fa+=1; }

// Region 13
// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr @stc, ptr [[BP0]]
// CK22-DAG: store ptr getelementptr inbounds nuw (i8, ptr @stc, i{{32|64}} 4),  ptr [[P0]]

// CK22: call void [[CALL13:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(stc [1:4])
  { stc[3].fa+=1; }

// Region 14

//  &std[0], &std[2], 5 * sizeof(std[0]), TO | FROM | ATTACH
//  &std,    &std[2], sizeof(std),        ATTACH

// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
// CK22-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
// CK22-DAG: [[RVAR0]] = load ptr, ptr @std
// CK22-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR00]] = load ptr, ptr @std

// CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
// CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CK22-DAG: store ptr @std, ptr [[BP1]]
// CK22-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
// CK22-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[RVAR1:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR1]] = load ptr, ptr @std

// CK22: call void [[CALL14:@.+]](ptr {{[^,]+}}, ptr null)
#pragma omp target map(std [2:5])
  { std[3].fa+=1; }

  return 0;
}
// CK22: define {{.+}}[[CALL00]]
// CK22: define {{.+}}[[CALL01]]
// CK22: define {{.+}}[[CALL02]]
// CK22: define {{.+}}[[CALL03]]
// CK22: define {{.+}}[[CALL04]]
// CK22: define {{.+}}[[CALL05]]
// CK22: define {{.+}}[[CALL06]]
// CK22: define {{.+}}[[CALL07]]
// CK22: define {{.+}}[[CALL08]]
// CK22: define {{.+}}[[CALL09]]
// CK22: define {{.+}}[[CALL10]]
// CK22: define {{.+}}[[CALL11]]
// CK22: define {{.+}}[[CALL12]]
// CK22: define {{.+}}[[CALL13]]
// CK22: define {{.+}}[[CALL14]]
#endif // CK22
#endif