1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4 
5 #ifdef CK1
6 
7 ///==========================================================================///
8 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -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 CK1
9 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
10 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK1
11 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -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 CK1
12 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
13 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK1
14 
15 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY10 %s
16 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
17 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY10 %s
18 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY10 %s
19 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
20 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY10 %s
21 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}}
22 
23 // CK1-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
24 
25 // CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
26 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
27 // CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
28 
29 // CK1-LABEL: implicit_maps_double_complex{{.*}}(
30 void implicit_maps_double_complex (int a){
31   double _Complex dc = (double)a;
32 
33   // CK1-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
34   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
35   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
36   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
37   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
38   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }**
39   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }**
40   // CK1-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]]
41   // CK1-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]]
42 
43   // CK1: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
44   #pragma omp target defaultmap(alloc:scalar)
45   {
46    dc *= dc;
47   }
48 }
49 
50 // CK1: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]])
51 // CK1: [[ADDR:%.+]] = alloca { double, double }*,
52 // CK1: store { double, double }* [[ARG]], { double, double }** [[ADDR]],
53 // CK1: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]],
54 // CK1: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0
55 #endif
56 ///==========================================================================///
57 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -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 CK2
58 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
59 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK2
60 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -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 CK2
61 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
62 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK2
63 
64 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY10 %s
65 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
66 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY10 %s
67 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY10 %s
68 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
69 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY10 %s
70 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}}
71 #ifdef CK2
72 
73 // CK2-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
74 
75 // CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
76 // Map types: OMP_MAP_TO  | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545
77 // CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545]
78 
79 // CK2-LABEL: implicit_maps_double_complex{{.*}}(
80 void implicit_maps_double_complex (int a){
81   double _Complex dc = (double)a;
82 
83   // CK2-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
84   // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
85   // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
86   // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
87   // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
88   // CK2-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }**
89   // CK2-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }**
90   // CK2-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]]
91   // CK2-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]]
92 
93   // CK2: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
94   #pragma omp target defaultmap(to:scalar)
95   {
96    dc *= dc;
97   }
98 }
99 
100 // CK2: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]])
101 // CK2: [[ADDR:%.+]] = alloca { double, double }*,
102 // CK2: store { double, double }* [[ARG]], { double, double }** [[ADDR]],
103 // CK2: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]],
104 // CK2: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0
105 #endif
106 ///==========================================================================///
107 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -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 CK3
108 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
109 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK3
110 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -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 CK3
111 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
112 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK3
113 
114 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY10 %s
115 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
116 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY10 %s
117 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY10 %s
118 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
119 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY10 %s
120 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}}
121 #ifdef CK3
122 
123 // CK3-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
124 
125 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
126 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546
127 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546]
128 
129 // CK3-LABEL: implicit_maps_double_complex{{.*}}(
130 void implicit_maps_double_complex (int a){
131   double _Complex dc = (double)a;
132 
133   // CK3-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
134   // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
135   // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
136   // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
137   // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
138   // CK3-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }**
139   // CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }**
140   // CK3-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]]
141   // CK3-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]]
142 
143   // CK3: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
144   #pragma omp target defaultmap(from:scalar)
145   {
146    dc *= dc;
147   }
148 }
149 
150 // CK3: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]])
151 // CK3: [[ADDR:%.+]] = alloca { double, double }*,
152 // CK3: store { double, double }* [[ARG]], { double, double }** [[ADDR]],
153 // CK3: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]],
154 // CK3: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0
155 #endif
156 ///==========================================================================///
157 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -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 CK4 --check-prefix CK4-64
158 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
159 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK4  --check-prefix CK4-64
160 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -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 CK4  --check-prefix CK4-32
161 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
162 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK4  --check-prefix CK4-32
163 
164 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY6 %s
165 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
166 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY6 %s
167 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY6 %s
168 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
169 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY6 %s
170 // SIMD-ONLY6-NOT: {{__kmpc|__tgt}}
171 #ifdef CK4
172 
173 // For a 32-bit targets, the value doesn't fit the size of the pointer,
174 // therefore it is passed by reference with a map 'to' specification.
175 
176 // CK4-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
177 
178 // CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
179 // Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
180 // CK4-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]
181 // Map types: OMP_MAP_TO  | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 673
182 // CK4-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673]
183 
184 // CK4-LABEL: implicit_maps_double{{.*}}(
185 void implicit_maps_double (int a){
186   double d = (double)a;
187 
188   // CK4-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
189   // CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
190   // CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
191   // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
192   // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
193 
194   // CK4-64-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz:64|32]]*
195   // CK4-64-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]*
196   // CK4-64-DAG: store i[[sz]] [[VAL:%[^,]+]], i[[sz]]* [[CBP1]]
197   // CK4-64-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]]
198   // CK4-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
199   // CK4-64-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to double*
200   // CK4-64-64-DAG: store double {{.+}}, double* [[CADDR]],
201 
202   // CK4-32-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**
203   // CK4-32-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
204   // CK4-32-DAG: store double* [[DECL:%[^,]+]], double** [[CBP1]]
205   // CK4-32-DAG: store double* [[DECL]], double** [[CP1]]
206 
207   // CK4-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
208   // CK4-32: call void [[KERNEL:@.+]](double* [[DECL]])
209   #pragma omp target defaultmap(firstprivate:scalar)
210   {
211     d += 1.0;
212   }
213 }
214 
215 // CK4-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
216 // CK4-64: [[ADDR:%.+]] = alloca i[[sz]],
217 // CK4-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
218 // CK4-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to double*
219 // CK4-64: {{.+}} = load double, double* [[CADDR]],
220 
221 // CK4-32: define internal void [[KERNEL]](double* {{.+}}[[ARG:%.+]])
222 // CK4-32: [[ADDR:%.+]] = alloca double*,
223 // CK4-32: store double* [[ARG]], double** [[ADDR]],
224 // CK4-32: [[REF:%.+]] = load double*, double** [[ADDR]],
225 // CK4-32: {{.+}} = load double, double* [[REF]],
226 
227 #endif
228 ///==========================================================================///
229 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -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 CK5
230 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
231 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK5
232 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -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 CK5
233 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
234 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK5
235 
236 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
237 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
238 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
239 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
240 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
241 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
242 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
243 #ifdef CK5
244 
245 // CK5-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
246 
247 // CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
248 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
249 // CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
250 
251 // CK5-LABEL: implicit_maps_array{{.*}}(
252 void implicit_maps_array (int a){
253   double darr[2] = {(double)a, (double)a};
254 
255   // CK5-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
256   // CK5-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
257   // CK5-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
258   // CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
259   // CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
260   // CK5-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]**
261   // CK5-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]**
262   // CK5-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]]
263   // CK5-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]]
264 
265   // CK5: call void [[KERNEL:@.+]]([2 x double]* [[DECL]])
266   #pragma omp target defaultmap(alloc: aggregate)
267   {
268     darr[0] += 1.0;
269     darr[1] += 1.0;
270   }
271 }
272 
273 // CK5: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]])
274 // CK5: [[ADDR:%.+]] = alloca [2 x double]*,
275 // CK5: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]],
276 // CK5: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]],
277 // CK5: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0
278 #endif
279 ///==========================================================================///
280 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -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 CK6
281 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
282 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK6
283 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -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 CK6
284 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
285 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK6
286 
287 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
288 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
289 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
290 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
291 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
292 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
293 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
294 #ifdef CK6
295 
296 // CK6-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
297 
298 // CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
299 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545
300 // CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545]
301 
302 // CK6-LABEL: implicit_maps_array{{.*}}(
303 void implicit_maps_array (int a){
304   double darr[2] = {(double)a, (double)a};
305 
306   // CK6-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
307   // CK6-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
308   // CK6-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
309   // CK6-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
310   // CK6-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
311   // CK6-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]**
312   // CK6-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]**
313   // CK6-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]]
314   // CK6-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]]
315 
316   // CK6: call void [[KERNEL:@.+]]([2 x double]* [[DECL]])
317 #pragma omp target defaultmap(to)
318   {
319     darr[0] += 1.0;
320     darr[1] += 1.0;
321   }
322 }
323 
324 // CK6: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]])
325 // CK6: [[ADDR:%.+]] = alloca [2 x double]*,
326 // CK6: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]],
327 // CK6: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]],
328 // CK6: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0
329 #endif
330 ///==========================================================================///
331 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -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 CK7
332 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
333 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK7
334 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -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 CK7
335 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
336 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK7
337 
338 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
339 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
340 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
341 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
342 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
343 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
344 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
345 #ifdef CK7
346 
347 // CK7-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
348 
349 // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
350 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546
351 // CK7-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546]
352 
353 // CK7-LABEL: implicit_maps_array{{.*}}(
354 void implicit_maps_array (int a){
355   double darr[2] = {(double)a, (double)a};
356 
357   // CK7-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
358   // CK7-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
359   // CK7-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
360   // CK7-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
361   // CK7-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
362   // CK7-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]**
363   // CK7-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]**
364   // CK7-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]]
365   // CK7-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]]
366 
367   // CK7: call void [[KERNEL:@.+]]([2 x double]* [[DECL]])
368   #pragma omp target defaultmap(from: aggregate)
369   {
370     darr[0] += 1.0;
371     darr[1] += 1.0;
372   }
373 }
374 
375 // CK7: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]])
376 // CK7: [[ADDR:%.+]] = alloca [2 x double]*,
377 // CK7: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]],
378 // CK7: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]],
379 // CK7: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0
380 #endif
381 ///==========================================================================///
382 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -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 CK8
383 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
384 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK8
385 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -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 CK8
386 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
387 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK8
388 
389 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
390 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
391 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
392 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
393 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
394 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
395 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
396 #ifdef CK8
397 
398 // CK8-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
399 
400 // CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
401 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
402 // CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]
403 
404 // CK8-LABEL: implicit_maps_array{{.*}}(
405 void implicit_maps_array (int a){
406   double darr[2] = {(double)a, (double)a};
407 
408   // CK8-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
409   // CK8-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
410   // CK8-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
411   // CK8-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
412   // CK8-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
413   // CK8-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]**
414   // CK8-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]**
415   // CK8-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]]
416   // CK8-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]]
417 
418   // CK8: call void [[KERNEL:@.+]]([2 x double]* [[DECL]])
419 #pragma omp target defaultmap(tofrom)
420   {
421     darr[0] += 1.0;
422     darr[1] += 1.0;
423   }
424 }
425 
426 // CK8: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]])
427 // CK8: [[ADDR:%.+]] = alloca [2 x double]*,
428 // CK8: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]],
429 // CK8: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]],
430 // CK8: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0
431 #endif
432 
433 ///==========================================================================///
434 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -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 CK9 --check-prefix CK9-64
435 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
436 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK9 --check-prefix CK9-64
437 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -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 CK9 --check-prefix CK9-32
438 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
439 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK9 --check-prefix CK9-32
440 
441 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY26 %s
442 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
443 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY26 %s
444 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY26 %s
445 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
446 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY26 %s
447 // SIMD-ONLY26-NOT: {{__kmpc|__tgt}}
448 
449 #ifdef CK9
450 
451 
452 // CK9-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
453 // CK9: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 40]
454 // CK9: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 673]
455 
456 
457 // CK9-LABEL: zero_size_section_and_private_maps{{.*}}(
458 void zero_size_section_and_private_maps (int ii){
459   int pvtArr[10];
460 
461   // Region 09
462   // CK9-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}, i8** null, i8** null)
463   // CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
464   // CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
465 
466   // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
467   // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
468   // CK9-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [10 x i32]**
469   // CK9-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [10 x i32]**
470   // CK9-DAG: store [10 x i32]* [[VAR0:%.+]], [10 x i32]** [[CBP0]]
471   // CK9-DAG: store [10 x i32]* [[VAR0]], [10 x i32]** [[CP0]]
472 
473   // CK9: call void [[CALL09:@.+]]([10 x i32]* {{[^,]+}})
474   #pragma omp target defaultmap(firstprivate:aggregate)
475   {
476     pvtArr[5]++;
477   }
478 
479 }
480 
481 
482 #endif
483 ///==========================================================================///
484 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -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 CK10
485 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
486 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK10
487 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -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 CK10
488 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
489 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK10
490 
491 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
492 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
493 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
494 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
495 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
496 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
497 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
498 #ifdef CK10
499 
500 
501 // CK10-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
502 
503 // CK10: [[SIZE:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
504 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
505 // CK10: [[MTYPE:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
506 
507 // CK10-LABEL: explicit_maps_single{{.*}}(
508 void explicit_maps_single (){
509   int *pa;
510 
511   // CK10-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}, i8** null, i8** null)
512   // CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
513   // CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
514 
515   // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
516   // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
517   // CK10-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
518   // CK10-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32***
519   // CK10-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
520   // CK10-DAG: store i32** [[VAR0]], i32*** [[CP0]]
521 
522   // CK10: call void [[CALL:@.+]](i32** {{[^,]+}})
523   #pragma omp target defaultmap(alloc: pointer)
524   {
525     pa[50]++;
526   }
527 }
528 
529   // CK10: define {{.+}}[[CALL]]
530 #endif
531 ///==========================================================================///
532 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -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 CK11
533 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
534 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK11
535 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -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 CK11
536 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
537 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK11
538 
539 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
540 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
541 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
542 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
543 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
544 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
545 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
546 #ifdef CK11
547 
548 
549 // CK11-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
550 
551 // CK11: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
552 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545
553 // CK11: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 545]
554 
555 // CK11-LABEL: explicit_maps_single{{.*}}(
556 void explicit_maps_single (){
557   int *pa;
558 
559   // CK11-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}, i8** null, i8** null)
560   // CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
561   // CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
562 
563   // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
564   // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
565   // CK11-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
566   // CK11-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32***
567   // CK11-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
568   // CK11-DAG: store i32** [[VAR0]], i32*** [[CP0]]
569 
570   // CK11: call void [[CALL09:@.+]](i32** {{[^,]+}})
571   #pragma omp target defaultmap(to: pointer)
572   {
573     pa[50]++;
574   }
575 }
576 
577   // CK11: define {{.+}}[[CALL09]]
578 #endif
579 ///==========================================================================///
580 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -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 CK12
581 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
582 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK12
583 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -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 CK12
584 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
585 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK12
586 
587 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
588 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
589 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
590 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
591 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
592 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
593 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
594 #ifdef CK12
595 
596 
597 // CK12-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
598 
599 // CK12: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
600 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546
601 // CK12: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 546]
602 
603 // CK12-LABEL: explicit_maps_single{{.*}}(
604 void explicit_maps_single (){
605   int *pa;
606 
607   // CK12-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}, i8** null, i8** null)
608   // CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
609   // CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
610 
611   // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
612   // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
613   // CK12-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
614   // CK12-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32***
615   // CK12-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
616   // CK12-DAG: store i32** [[VAR0]], i32*** [[CP0]]
617 
618   // CK12: call void [[CALL09:@.+]](i32** {{[^,]+}})
619   #pragma omp target defaultmap(from: pointer)
620   {
621     pa[50]++;
622   }
623 }
624 
625   // CK12: define {{.+}}[[CALL09]]
626 #endif
627 ///==========================================================================///
628 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -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 CK13
629 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
630 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK13
631 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -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 CK13
632 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
633 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK13
634 
635 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
636 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
637 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
638 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
639 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
640 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
641 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
642 #ifdef CK13
643 
644 
645 // CK13-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
646 
647 // CK13: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
648 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
649 // CK13: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 547]
650 
651 // CK13-LABEL: explicit_maps_single{{.*}}(
652 void explicit_maps_single (){
653   int *pa;
654 
655   // CK13-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}, i8** null, i8** null)
656   // CK13-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
657   // CK13-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
658 
659   // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
660   // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
661   // CK13-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
662   // CK13-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32***
663   // CK13-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
664   // CK13-DAG: store i32** [[VAR0]], i32*** [[CP0]]
665 
666   // CK13: call void [[CALL09:@.+]](i32** {{[^,]+}})
667   #pragma omp target defaultmap(tofrom: pointer)
668   {
669     pa[50]++;
670   }
671 }
672 
673   // CK13: define {{.+}}[[CALL09]]
674 #endif
675 ///==========================================================================///
676 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -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 CK14
677 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
678 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK14
679 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -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 CK14
680 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
681 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK14
682 
683 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
684 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
685 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
686 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY8 %s
687 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
688 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY8 %s
689 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}}
690 #ifdef CK14
691 
692 
693 // CK14-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
694 
695 // CK14: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
696 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
697 // CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
698 
699 // CK14-LABEL: explicit_maps_single{{.*}}(
700 void explicit_maps_single (){
701   int *pa;
702 
703   // CK14-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}, i8** null, i8** null)
704   // CK14-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
705   // CK14-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
706 
707   // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
708   // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
709   // CK14-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
710   // CK14-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
711   // CK14-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
712   // CK14-DAG: store i32* [[VAR0]], i32** [[CP0]]
713 
714   // CK14: call void [[CALL09:@.+]](i32* {{[^,]+}})
715   #pragma omp target defaultmap(firstprivate: pointer)
716   {
717     pa[50]++;
718   }
719 }
720 
721   // CK14: define {{.+}}[[CALL09]]
722 #endif
723 ///==========================================================================///
724 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -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 CK15
725 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
726 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK15
727 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -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 CK15
728 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
729 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK15
730 
731 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY12 %s
732 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
733 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY12 %s
734 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY12 %s
735 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
736 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY12 %s
737 // SIMD-ONLY12-NOT: {{__kmpc|__tgt}}
738 #ifdef CK15
739 
740 // CK15-LABEL: @.__omp_offloading_{{.*}}implicit_maps_variable_length_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
741 
742 // CK15-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 {{4|8}}, i64 {{4|8}}, i64 0]
743 // Map types:
744 //  - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size)
745 //  - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size)
746 //  - OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
747 // CK15-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 800, i64 800, i64 544]
748 
749 // CK15-LABEL: implicit_maps_variable_length_array{{.*}}(
750 void implicit_maps_variable_length_array (int a){
751   double vla[2][a];
752 
753   // CK15-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], i64* [[SGEP:%[^,]+]], {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
754   // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
755   // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
756   // CK15-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SS:%[^,]+]], i32 0, i32 0
757 
758   // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
759   // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
760   // CK15-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[sz:64|32]]*
761   // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[sz]]*
762   // CK15-DAG: store i[[sz]] 2, i[[sz]]* [[CBP0]]
763   // CK15-DAG: store i[[sz]] 2, i[[sz]]* [[CP0]]
764 
765   // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
766   // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
767   // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]*
768   // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]*
769   // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]]
770   // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]]
771 
772   // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
773   // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
774   // CK15-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 2
775   // CK15-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double**
776   // CK15-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double**
777   // CK15-DAG: store double* [[DECL:%.+]], double** [[CBP2]]
778   // CK15-DAG: store double* [[DECL]], double** [[CP2]]
779   // CK15-DAG: store i64 [[VALS2:%.+]], i64* [[S2]],
780   // CK15-DAG: [[VALS2]] = {{mul nuw i64 %.+, 8|sext i32 %.+ to i64}}
781 
782   // CK15: call void [[KERNEL:@.+]](i[[sz]] {{.+}}, i[[sz]] {{.+}}, double* [[DECL]])
783   #pragma omp target defaultmap(alloc: aggregate)
784   {
785     vla[1][3] += 1.0;
786   }
787 }
788 
789 // CK15: define internal void [[KERNEL]](i[[sz]] [[VLA0:%.+]], i[[sz]] [[VLA1:%.+]], double* {{.*}}[[ARG:%.+]])
790 // CK15: [[ADDR0:%.+]] = alloca i[[sz]],
791 // CK15: [[ADDR1:%.+]] = alloca i[[sz]],
792 // CK15: [[ADDR2:%.+]] = alloca double*,
793 // CK15: store i[[sz]] [[VLA0]], i[[sz]]* [[ADDR0]],
794 // CK15: store i[[sz]] [[VLA1]], i[[sz]]* [[ADDR1]],
795 // CK15: store double* [[ARG]], double** [[ADDR2]],
796 // CK15: {{.+}} = load i[[sz]],  i[[sz]]* [[ADDR0]],
797 // CK15: {{.+}} = load i[[sz]],  i[[sz]]* [[ADDR1]],
798 // CK15: [[REF:%.+]] = load double*, double** [[ADDR2]],
799 // CK15: {{.+}} = getelementptr inbounds double, double* [[REF]], i[[sz]] %{{.+}}
800 #endif
801 ///==========================================================================///
802 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -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 CK16
803 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
804 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK16
805 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -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 CK16
806 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
807 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK16
808 
809 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
810 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
811 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
812 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
813 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
814 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
815 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}}
816 #ifdef CK16
817 
818 // CK16-DAG: [[ST:%.+]] = type { i32, double }
819 // CK16-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
820 // CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]
821 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
822 // CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
823 
824 class SSS {
825 public:
826   int a;
827   double b;
828 };
829 
830 // CK16-LABEL: implicit_maps_struct{{.*}}(
831 void implicit_maps_struct (int a){
832   SSS s = {a, (double)a};
833 
834   // CK16-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
835   // CK16-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
836   // CK16-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
837   // CK16-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
838   // CK16-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
839   // CK16-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
840   // CK16-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]**
841   // CK16-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]]
842   // CK16-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]]
843 
844   // CK16: call void [[KERNEL:@.+]]([[ST]]* [[DECL]])
845   #pragma omp target defaultmap(alloc: aggregate)
846   {
847     s.a += 1;
848     s.b += 1.0;
849   }
850 }
851 
852 // CK16: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]])
853 // CK16: [[ADDR:%.+]] = alloca [[ST]]*,
854 // CK16: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]],
855 // CK16: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]],
856 // CK16: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0
857 #endif
858 ///==========================================================================///
859 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -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 CK17
860 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
861 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK17
862 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -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 CK17
863 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
864 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK17
865 
866 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
867 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
868 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
869 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
870 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
871 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
872 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}}
873 #ifdef CK17
874 
875 // CK17-DAG: [[ST:%.+]] = type { i32, double }
876 // CK17-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
877 // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]
878 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545
879 // CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545]
880 
881 class SSS {
882 public:
883   int a;
884   double b;
885 };
886 
887 // CK17-LABEL: implicit_maps_struct{{.*}}(
888 void implicit_maps_struct (int a){
889   SSS s = {a, (double)a};
890 
891   // CK17-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
892   // CK17-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
893   // CK17-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
894   // CK17-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
895   // CK17-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
896   // CK17-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
897   // CK17-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]**
898   // CK17-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]]
899   // CK17-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]]
900 
901   // CK17: call void [[KERNEL:@.+]]([[ST]]* [[DECL]])
902   #pragma omp target defaultmap(to: aggregate)
903   {
904     s.a += 1;
905     s.b += 1.0;
906   }
907 }
908 
909 // CK17: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]])
910 // CK17: [[ADDR:%.+]] = alloca [[ST]]*,
911 // CK17: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]],
912 // CK17: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]],
913 // CK17: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0
914 #endif
915 ///==========================================================================///
916 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -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 CK18
917 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
918 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK18
919 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -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 CK18
920 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
921 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK18
922 
923 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
924 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
925 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
926 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
927 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
928 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
929 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}}
930 #ifdef CK18
931 
932 // CK18-DAG: [[ST:%.+]] = type { i32, double }
933 // CK18-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
934 // CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]
935 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546
936 // CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546]
937 
938 class SSS {
939 public:
940   int a;
941   double b;
942 };
943 
944 // CK18-LABEL: implicit_maps_struct{{.*}}(
945 void implicit_maps_struct (int a){
946   SSS s = {a, (double)a};
947 
948   // CK18-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
949   // CK18-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
950   // CK18-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
951   // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
952   // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
953   // CK18-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
954   // CK18-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]**
955   // CK18-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]]
956   // CK18-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]]
957 
958   // CK18: call void [[KERNEL:@.+]]([[ST]]* [[DECL]])
959   #pragma omp target defaultmap(from: aggregate)
960   {
961     s.a += 1;
962     s.b += 1.0;
963   }
964 }
965 
966 // CK18: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]])
967 // CK18: [[ADDR:%.+]] = alloca [[ST]]*,
968 // CK18: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]],
969 // CK18: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]],
970 // CK18: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0
971 #endif
972 ///==========================================================================///
973 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -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 CK19
974 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
975 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK19
976 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -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 CK19
977 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
978 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK19
979 
980 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
981 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
982 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
983 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
984 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
985 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
986 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}}
987 #ifdef CK19
988 
989 // CK19-DAG: [[ST:%.+]] = type { i32, double }
990 // CK19-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
991 // CK19-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]
992 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
993 // CK19-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]
994 
995 class SSS {
996 public:
997   int a;
998   double b;
999 };
1000 
1001 // CK19-LABEL: implicit_maps_struct{{.*}}(
1002 void implicit_maps_struct (int a){
1003   SSS s = {a, (double)a};
1004 
1005   // CK19-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
1006   // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1007   // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1008   // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1009   // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1010   // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
1011   // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]**
1012   // CK19-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]]
1013   // CK19-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]]
1014 
1015   // CK19: call void [[KERNEL:@.+]]([[ST]]* [[DECL]])
1016   #pragma omp target defaultmap(tofrom: aggregate)
1017   {
1018     s.a += 1;
1019     s.b += 1.0;
1020   }
1021 }
1022 
1023 // CK19: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]])
1024 // CK19: [[ADDR:%.+]] = alloca [[ST]]*,
1025 // CK19: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]],
1026 // CK19: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]],
1027 // CK19: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0
1028 #endif
1029 ///==========================================================================///
1030 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -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 CK20 --check-prefix CK20-64
1031 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1032 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK20  --check-prefix CK20-64
1033 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -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 CK20  --check-prefix CK20-32
1034 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1035 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK20  --check-prefix CK20-32
1036 
1037 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY6 %s
1038 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1039 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY6 %s
1040 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY6 %s
1041 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1042 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY6 %s
1043 // SIMD-ONLY6-NOT: {{__kmpc|__tgt}}
1044 #ifdef CK20
1045 
1046 // For a 32-bit targets, the value doesn't fit the size of the pointer,
1047 // therefore it is passed by reference with a map 'to' specification.
1048 
1049 // CK20-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
1050 
1051 // CK20-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
1052 // Map types: OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
1053 // CK20-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]
1054 // Map types: OMP_MAP_TO | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 673
1055 // CK20-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673]
1056 
1057 // CK20-LABEL: implicit_maps_double{{.*}}(
1058 void implicit_maps_double (int a){
1059   double d = (double)a;
1060 
1061   // CK20-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
1062   // CK20-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1063   // CK20-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1064   // CK20-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1065   // CK20-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1066 
1067   // CK20-64-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz:64|32]]*
1068   // CK20-64-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]*
1069   // CK20-64-DAG: store i[[sz]] [[VAL:%[^,]+]], i[[sz]]* [[CBP1]]
1070   // CK20-64-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]]
1071   // CK20-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
1072   // CK20-64-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to double*
1073   // CK20-64-64-DAG: store double {{.+}}, double* [[CADDR]],
1074 
1075   // CK20-32-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**
1076   // CK20-32-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
1077   // CK20-32-DAG: store double* [[DECL:%[^,]+]], double** [[CBP1]]
1078   // CK20-32-DAG: store double* [[DECL]], double** [[CP1]]
1079 
1080   // CK20-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
1081   // CK20-32: call void [[KERNEL:@.+]](double* [[DECL]])
1082   #pragma omp target defaultmap(default: scalar)
1083   {
1084     d += 1.0;
1085   }
1086 }
1087 
1088 // CK20-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
1089 // CK20-64: [[ADDR:%.+]] = alloca i[[sz]],
1090 // CK20-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
1091 // CK20-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to double*
1092 // CK20-64: {{.+}} = load double, double* [[CADDR]],
1093 
1094 // CK20-32: define internal void [[KERNEL]](double* {{.+}}[[ARG:%.+]])
1095 // CK20-32: [[ADDR:%.+]] = alloca double*,
1096 // CK20-32: store double* [[ARG]], double** [[ADDR]],
1097 // CK20-32: [[REF:%.+]] = load double*, double** [[ADDR]],
1098 // CK20-32: {{.+}} = load double, double* [[REF]],
1099 
1100 #endif
1101 ///==========================================================================///
1102 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -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 CK21
1103 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1104 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK21
1105 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -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 CK21
1106 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1107 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK21
1108 
1109 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
1110 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1111 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
1112 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY16 %s
1113 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1114 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY16 %s
1115 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}}
1116 #ifdef CK21
1117 
1118 // CK21-DAG: [[ST:%.+]] = type { i32, double }
1119 // CK21-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
1120 // CK21-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]
1121 // Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
1122 // CK21-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]
1123 
1124 class SSS {
1125 public:
1126   int a;
1127   double b;
1128 };
1129 
1130 // CK21-LABEL: implicit_maps_struct{{.*}}(
1131 void implicit_maps_struct (int a){
1132   SSS s = {a, (double)a};
1133 
1134   // CK21-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
1135   // CK21-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1136   // CK21-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1137   // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1138   // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1139   // CK21-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
1140   // CK21-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]**
1141   // CK21-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]]
1142   // CK21-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]]
1143 
1144   // CK21: call void [[KERNEL:@.+]]([[ST]]* [[DECL]])
1145   #pragma omp target defaultmap(default: aggregate)
1146   {
1147     s.a += 1;
1148     s.b += 1.0;
1149   }
1150 }
1151 
1152 // CK21: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]])
1153 // CK21: [[ADDR:%.+]] = alloca [[ST]]*,
1154 // CK21: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]],
1155 // CK21: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]],
1156 // CK21: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0
1157 #endif
1158 ///==========================================================================///
1159 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1160 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1161 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1162 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1163 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1164 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1165 
1166 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY9 %s
1167 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1168 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY9 %s
1169 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY9 %s
1170 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1171 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY9 %s
1172 // SIMD-ONLY9-NOT: {{__kmpc|__tgt}}
1173 #ifdef CK22
1174 
1175 // CK22-LABEL: @.__omp_offloading_{{.*}}implicit_maps_pointer{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
1176 
1177 // CK22-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
1178 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544
1179 // CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]
1180 
1181 // CK22-LABEL: implicit_maps_pointer{{.*}}(
1182 void implicit_maps_pointer (){
1183   double *ddyn;
1184 
1185   // CK22-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
1186   // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1187   // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1188   // CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1189   // CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1190   // CK22-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**
1191   // CK22-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
1192   // CK22-DAG: store double* [[PTR:%[^,]+]], double** [[CBP1]]
1193   // CK22-DAG: store double* [[PTR]], double** [[CP1]]
1194 
1195   // CK22: call void [[KERNEL:@.+]](double* [[PTR]])
1196   #pragma omp target defaultmap(default: pointer)
1197   {
1198     ddyn[0] += 1.0;
1199     ddyn[1] += 1.0;
1200   }
1201 }
1202 
1203 // CK22: define internal void [[KERNEL]](double* {{.*}}[[ARG:%.+]])
1204 // CK22: [[ADDR:%.+]] = alloca double*,
1205 // CK22: store double* [[ARG]], double** [[ADDR]],
1206 // CK22: [[REF:%.+]] = load double*, double** [[ADDR]],
1207 // CK22: {{.+}} = getelementptr inbounds double, double* [[REF]], i{{64|32}} 0
1208 
1209 #endif
1210 ///==========================================================================///
1211 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64
1212 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1213 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 %s  --check-prefix CK23 --check-prefix CK23-64
1214 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK23 --check-prefix CK23-32
1215 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1216 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 %s  --check-prefix CK23 --check-prefix CK23-32
1217 
1218 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
1219 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1220 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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 --check-prefix SIMD-ONLY0 %s
1221 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
1222 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1223 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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 --check-prefix SIMD-ONLY0 %s
1224 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
1225 #ifdef CK23
1226 
1227 double *g;
1228 
1229 // CK23: @g ={{.*}} global double*
1230 // CK23: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
1231 // CK23: [[TYPES00:@.+]] = {{.+}}constant [1 x i64] [i64 288]
1232 
1233 // CK23: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
1234 // CK23: [[TYPES01:@.+]] = {{.+}}constant [1 x i64] [i64 288]
1235 
1236 // CK23: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
1237 // CK23: [[TYPES02:@.+]] = {{.+}}constant [1 x i64] [i64 288]
1238 
1239 // CK23: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
1240 // CK23: [[TYPES03:@.+]] = {{.+}}constant [1 x i64] [i64 288]
1241 
1242 // CK23: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
1243 // CK23: [[TYPES04:@.+]] = {{.+}}constant [1 x i64] [i64 288]
1244 
1245 // CK23: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
1246 // CK23: [[TYPES05:@.+]] = {{.+}}constant [1 x i64] [i64 288]
1247 
1248 // CK23: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
1249 // CK23: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288]
1250 
1251 // CK23-LABEL: @_Z3foo{{.*}}(
1252 template<typename T>
1253 void foo(float *&lr, T *&tr) {
1254   float *l;
1255   T *t;
1256 
1257   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}}, i8** null, i8** null)
1258   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1259   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1260   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1261   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1262   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**
1263   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
1264   // CK23-DAG: store double* [[VAL:%.+]], double** [[CBP1]]
1265   // CK23-DAG: store double* [[VAL]], double** [[CP1]]
1266   // CK23-DAG: [[VAL]] = load double*, double** [[ADDR:@g]],
1267 
1268   // CK23: call void [[KERNEL:@.+]](double* [[VAL]])
1269   #pragma omp target is_device_ptr(g) defaultmap(none:pointer)
1270   {
1271     ++g;
1272   }
1273 
1274   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}}, i8** null, i8** null)
1275   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1276   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1277   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1278   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1279   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
1280   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
1281   // CK23-DAG: store float* [[VAL:%.+]], float** [[CBP1]]
1282   // CK23-DAG: store float* [[VAL]], float** [[CP1]]
1283   // CK23-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
1284 
1285   // CK23: call void [[KERNEL:@.+]](float* [[VAL]])
1286   #pragma omp target is_device_ptr(l) defaultmap(none:pointer)
1287   {
1288     ++l;
1289   }
1290 
1291   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}}, i8** null, i8** null)
1292   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1293   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1294   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1295   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1296   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
1297   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
1298   // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
1299   // CK23-DAG: store i32* [[VAL]], i32** [[CP1]]
1300   // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
1301 
1302   // CK23: call void [[KERNEL:@.+]](i32* [[VAL]])
1303   #pragma omp target is_device_ptr(t) defaultmap(none:pointer)
1304   {
1305     ++t;
1306   }
1307 
1308   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}}, i8** null, i8** null)
1309   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1310   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1311   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1312   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1313   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
1314   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
1315   // CK23-DAG: store float* [[VAL:%.+]], float** [[CBP1]]
1316   // CK23-DAG: store float* [[VAL]], float** [[CP1]]
1317   // CK23-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
1318   // CK23-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]],
1319 
1320   // CK23: call void [[KERNEL:@.+]](float* [[VAL]])
1321   #pragma omp target is_device_ptr(lr) defaultmap(none:pointer)
1322   {
1323     ++lr;
1324   }
1325 
1326   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}}, i8** null, i8** null)
1327   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1328   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1329   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1330   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1331   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
1332   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
1333   // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
1334   // CK23-DAG: store i32* [[VAL]], i32** [[CP1]]
1335   // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
1336   // CK23-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
1337 
1338   // CK23: call void [[KERNEL:@.+]](i32* [[VAL]])
1339   #pragma omp target is_device_ptr(tr) defaultmap(none:pointer)
1340   {
1341     ++tr;
1342   }
1343 
1344   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}}, i8** null, i8** null)
1345   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1346   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1347   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1348   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1349   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
1350   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
1351   // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
1352   // CK23-DAG: store i32* [[VAL]], i32** [[CP1]]
1353   // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
1354   // CK23-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
1355 
1356   // CK23: call void [[KERNEL:@.+]](i32* [[VAL]])
1357   #pragma omp target is_device_ptr(tr,lr) defaultmap(none:pointer)
1358   {
1359     ++tr;
1360   }
1361 
1362   // CK23-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}}, i8** null, i8** null)
1363   // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1364   // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1365   // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
1366   // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
1367   // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
1368   // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
1369   // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
1370   // CK23-DAG: store i32* [[VAL]], i32** [[CP1]]
1371   // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
1372   // CK23-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
1373 
1374   // CK23-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
1375   // CK23-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
1376   // CK23-DAG: [[_CBP1:%.+]] = bitcast i8** [[_BP1]] to float**
1377   // CK23-DAG: [[_CP1:%.+]] = bitcast i8** [[_P1]] to float**
1378   // CK23-DAG: store float* [[_VAL:%.+]], float** [[_CBP1]]
1379   // CK23-DAG: store float* [[_VAL]], float** [[_CP1]]
1380   // CK23-DAG: [[_VAL]] = load float*, float** [[_ADDR:%.+]],
1381   // CK23-DAG: [[_ADDR]] = load float**, float*** [[_ADDR2:%.+]],
1382 
1383   // CK23: call void [[KERNEL:@.+]](i32* [[VAL]], float* [[_VAL]])
1384   #pragma omp target is_device_ptr(tr,lr) defaultmap(none:pointer)
1385   {
1386     ++tr,++lr;
1387   }
1388 }
1389 
1390 void bar(float *&a, int *&b) {
1391   foo<int>(a,b);
1392 }
1393 
1394 #endif
1395 ///==========================================================================///
1396 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1397 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1398 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1399 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1400 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1401 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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
1402 
1403 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
1404 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1405 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY18 %s
1406 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
1407 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1408 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY18 %s
1409 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
1410 #ifdef CK24
1411 
1412 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
1413 // CK24: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
1414 // CK24: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059]
1415 
1416 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
1417 // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
1418 // CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063]
1419 
1420 // CK24-LABEL: explicit_maps_single{{.*}}(
1421 void explicit_maps_single (int ii){
1422   // Map of a scalar.
1423   int a = ii;
1424 
1425   // Close.
1426   // Region 00
1427   // CK24-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null)
1428   // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1429   // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1430 
1431   // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1432   // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1433   // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
1434   // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
1435   // CK24-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
1436   // CK24-DAG: store i32* [[VAR0]], i32** [[CP0]]
1437 
1438   // CK24: call void [[CALL00:@.+]](i32* {{[^,]+}})
1439   #pragma omp target map(close, tofrom: a) defaultmap(none:scalar)
1440   {
1441    a++;
1442   }
1443 
1444   // Always Close.
1445   // Region 01
1446   // CK24-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null, i8** null)
1447   // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1448   // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1449 
1450   // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1451   // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1452   // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
1453   // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
1454   // CK24-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
1455   // CK24-DAG: store i32* [[VAR0]], i32** [[CP0]]
1456 
1457   // CK24: call void [[CALL01:@.+]](i32* {{[^,]+}})
1458   #pragma omp target map(always close tofrom: a) defaultmap(none:scalar)
1459   {
1460    a++;
1461   }
1462 }
1463 // CK24: define {{.+}}[[CALL00]]
1464 // CK24: define {{.+}}[[CALL01]]
1465 
1466 #endif
1467 ///==========================================================================///
1468 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -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 CK25 --check-prefix CK25-64
1469 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1470 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK25 --check-prefix CK25-64
1471 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -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 CK25 --check-prefix CK25-32
1472 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1473 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK25 --check-prefix CK25-32
1474 
1475 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
1476 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1477 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY18 %s
1478 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
1479 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1480 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY18 %s
1481 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
1482 #ifdef CK25
1483 
1484 extern int x;
1485 #pragma omp declare target to(x)
1486 
1487 // CK25-LABEL: @.__omp_offloading_{{.*}}declare_target_to{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0
1488 
1489 void declare_target_to()
1490 {
1491   // CK25: [[C:%.+]] = load i32, i32* @x
1492   // CK25: [[INC:%.+]] = add nsw i32 [[C]], 1
1493   // CK25: store i32 [[INC]], i32* @x
1494   // CK25: ret void
1495   #pragma omp target defaultmap(none : scalar)
1496   {
1497     x++;
1498   }
1499 }
1500 
1501 #endif
1502 ///==========================================================================///
1503 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -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 CK26 --check-prefix CK26-64
1504 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1505 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK26 --check-prefix CK26-64
1506 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -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 CK26 --check-prefix CK26-32
1507 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1508 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -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 CK26 --check-prefix CK26-32
1509 
1510 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
1511 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
1512 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY18 %s
1513 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
1514 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
1515 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -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  --check-prefix SIMD-ONLY18 %s
1516 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
1517 #ifdef CK26
1518 
1519 // CK26-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4096, i64 {{.+}}]
1520 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_PTR_AND_OBJ | OMP_MAP_IMPLICIT = 531
1521 // CK26-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 531, i64 531, i64 531]
1522 
1523 float Vector[1024];
1524 #pragma omp declare target link(Vector)
1525 
1526 extern int a;
1527 #pragma omp declare target link(a)
1528 
1529 double *ptr;
1530 #pragma omp declare target link(ptr)
1531 
1532 void declare_target_link()
1533 {
1534 #pragma omp target defaultmap(none:scalar) defaultmap(none:aggregate) defaultmap(none:pointer)
1535   {
1536 
1537     // CK26-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null)
1538     // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
1539     // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
1540     // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
1541     // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
1542     // CK26-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double****
1543     // CK26-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
1544     // CK26-DAG: store double*** @ptr_decl_tgt_ref_ptr, double**** [[CBP1]]
1545     // CK26-DAG: store double** @ptr, double*** [[CP1]]
1546 
1547     Vector[a]++;
1548     ptr++;
1549   }
1550 }
1551 
1552 #endif
1553 #endif
1554