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