1 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
2 // RUN: %clang_cc1 -no-opaque-pointers -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
3 // RUN: %clang_cc1 -no-opaque-pointers -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
4
5 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -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
6 // RUN: %clang_cc1 -no-opaque-pointers -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
7 // RUN: %clang_cc1 -no-opaque-pointers -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
8 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
9 // expected-no-diagnostics
10
11 #ifndef HEADER
12 #define HEADER
13
14 // CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
15 // 64 = 0x40 = OMP_MAP_RETURN_PARAM
16 // CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 64, i64 64, i64 64, i64 64, i64 64]
17 // CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
18 // 0 = OMP_MAP_NONE
19 // 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM
20 // CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 0, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720]
21 struct S {
22 int a = 0;
23 int *ptr = &a;
24 int &ref = a;
25 int arr[4];
SS26 S() {}
fooS27 void foo() {
28 #pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a])
29 ++a, ++*ptr, ++ref, ++arr[0];
30 }
31 };
32
main()33 int main() {
34 float a = 0;
35 float *ptr = &a;
36 float &ref = a;
37 float arr[4];
38 float vla[(int)a];
39 S s;
40 s.foo();
41 #pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0])
42 ++a, ++*ptr, ++ref, ++arr[0], ++vla[0];
43 return a;
44 }
45
46 // CHECK-LABEL: @main()
47 // CHECK: [[A_ADDR:%.+]] = alloca float,
48 // CHECK: [[PTR_ADDR:%.+]] = alloca float*,
49 // CHECK: [[REF_ADDR:%.+]] = alloca float*,
50 // CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float],
51 // CHECK: [[BPTRS:%.+]] = alloca [5 x i8*],
52 // CHECK: [[PTRS:%.+]] = alloca [5 x i8*],
53 // CHECK: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}},
54 // CHECK: [[PTR:%.+]] = load float*, float** [[PTR_ADDR]],
55 // CHECK: [[REF:%.+]] = load float*, float** [[REF_ADDR]],
56 // CHECK: [[ARR:%.+]] = getelementptr inbounds [4 x float], [4 x float]* [[ARR_ADDR]], i64 0, i64 0
57 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
58 // CHECK: [[BPTR0_A_ADDR:%.+]] = bitcast i8** [[BPTR0]] to float**
59 // CHECK: store float* [[A_ADDR]], float** [[BPTR0_A_ADDR]],
60 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
61 // CHECK: [[PTR0_A_ADDR:%.+]] = bitcast i8** [[PTR0]] to float**
62 // CHECK: store float* [[A_ADDR]], float** [[PTR0_A_ADDR]],
63 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 1
64 // CHECK: [[BPTR1_PTR_ADDR:%.+]] = bitcast i8** [[BPTR1]] to float**
65 // CHECK: store float* [[PTR]], float** [[BPTR1_PTR_ADDR]],
66 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 1
67 // CHECK: [[PTR1_PTR_ADDR:%.+]] = bitcast i8** [[PTR1]] to float**
68 // CHECK: store float* [[PTR]], float** [[PTR1_PTR_ADDR]],
69 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 2
70 // CHECK: [[BPTR2_REF_ADDR:%.+]] = bitcast i8** [[BPTR2]] to float**
71 // CHECK: store float* [[REF]], float** [[BPTR2_REF_ADDR]],
72 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 2
73 // CHECK: [[PTR2_REF_ADDR:%.+]] = bitcast i8** [[PTR2]] to float**
74 // CHECK: store float* [[REF]], float** [[PTR2_REF_ADDR]],
75 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 3
76 // CHECK: [[BPTR3_ARR_ADDR:%.+]] = bitcast i8** [[BPTR3]] to float**
77 // CHECK: store float* [[ARR]], float** [[BPTR3_ARR_ADDR]],
78 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 3
79 // CHECK: [[PTR3_ARR_ADDR:%.+]] = bitcast i8** [[PTR3]] to float**
80 // CHECK: store float* [[ARR]], float** [[PTR3_ARR_ADDR]],
81 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 4
82 // CHECK: [[BPTR4_VLA_ADDR:%.+]] = bitcast i8** [[BPTR4]] to float**
83 // CHECK: store float* [[VLA_ADDR]], float** [[BPTR4_VLA_ADDR]],
84 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 4
85 // CHECK: [[PTR4_VLA_ADDR:%.+]] = bitcast i8** [[PTR4]] to float**
86 // CHECK: store float* [[VLA_ADDR]], float** [[PTR4_VLA_ADDR]],
87 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
88 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
89 // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[SIZES1]], i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null)
90 // CHECK: [[A_REF:%.+]] = load float*, float** [[BPTR0_A_ADDR]],
91 // CHECK: [[REF_REF:%.+]] = load float*, float** [[BPTR2_REF_ADDR]],
92 // CHECK: store float* [[REF_REF]], float** [[TMP_REF_ADDR:%.+]],
93 // CHECK: [[BPTR3_ARR_ADDR_CAST:%.+]] = bitcast float** [[BPTR3_ARR_ADDR]] to [4 x float]**
94 // CHECK: [[ARR_REF:%.+]] = load [4 x float]*, [4 x float]** [[BPTR3_ARR_ADDR_CAST]],
95 // CHECK: [[VLA_REF:%.+]] = load float*, float** [[BPTR4_VLA_ADDR]],
96 // CHECK: [[A:%.+]] = load float, float* [[A_REF]],
97 // CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00
98 // CHECK: store float [[INC]], float* [[A_REF]],
99 // CHECK: [[PTR_ADDR:%.+]] = load float*, float** [[BPTR1_PTR_ADDR]],
100 // CHECK: [[VAL:%.+]] = load float, float* [[PTR_ADDR]],
101 // CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00
102 // CHECK: store float [[INC]], float* [[PTR_ADDR]],
103 // CHECK: [[REF_ADDR:%.+]] = load float*, float** [[TMP_REF_ADDR]],
104 // CHECK: [[REF:%.+]] = load float, float* [[REF_ADDR]],
105 // CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00
106 // CHECK: store float [[INC]], float* [[REF_ADDR]],
107 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x float], [4 x float]* [[ARR_REF]], i64 0, i64 0
108 // CHECK: [[ARR0:%.+]] = load float, float* [[ARR0_ADDR]],
109 // CHECK: [[INC:%.+]] = fadd float [[ARR0]], 1.000000e+00
110 // CHECK: store float [[INC]], float* [[ARR0_ADDR]],
111 // CHECK: [[VLA0_ADDR:%.+]] = getelementptr inbounds float, float* [[VLA_REF]], i64 0
112 // CHECK: [[VLA0:%.+]] = load float, float* [[VLA0_ADDR]],
113 // CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00
114 // CHECK: store float [[INC]], float* [[VLA0_ADDR]],
115 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
116 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
117 // CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[SIZES1]], i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null)
118
119 // CHECK: foo
120 // %this.addr = alloca %struct.S*, align 8
121 // CHECK: [[BPTRS:%.+]] = alloca [5 x i8*],
122 // CHECK: [[PTRS:%.+]] = alloca [5 x i8*],
123 // CHECK: [[SIZES:%.+]] = alloca [5 x i64],
124 // %tmp = alloca i32*, align 8
125 // %tmp6 = alloca i32**, align 8
126 // %tmp7 = alloca i32*, align 8
127 // %tmp8 = alloca i32**, align 8
128 // %tmp9 = alloca [4 x i32]*, align 8
129 // store %struct.S* %this, %struct.S** %this.addr, align 8
130 // %this1 = load %struct.S*, %struct.S** %this.addr, align 8
131 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS:%.+]], i32 0, i32 0
132 // %ptr = getelementptr inbounds %struct.S, %struct.S* %this1, i32 0, i32 1
133 // %ref = getelementptr inbounds %struct.S, %struct.S* %this1, i32 0, i32 2
134 // %0 = load i32*, i32** %ref, align 8
135 // CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3
136 // CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 0
137 // CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 1
138 // CHECK: [[REF_REF:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 2
139 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[REF_REF]],
140 // CHECK: [[ARR_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3
141 // CHECK: [[ARR_END:%.+]] = getelementptr [4 x i32], [4 x i32]* [[ARR_ADDR]], i32 1
142 // CHECK: [[BEGIN:%.+]] = bitcast i32* [[A_ADDR]] to i8*
143 // CHECK: [[END:%.+]] = bitcast [4 x i32]* [[ARR_END]] to i8*
144 // CHECK: [[E:%.+]] = ptrtoint i8* [[END]] to i64
145 // CHECK: [[B:%.+]] = ptrtoint i8* [[BEGIN]] to i64
146 // CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]]
147 // CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
148 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
149 // CHECK: [[BPTR0_S:%.+]] = bitcast i8** [[BPTR0]] to %struct.S**
150 // CHECK: store %struct.S* [[THIS]], %struct.S** [[BPTR0_S]],
151 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
152 // CHECK: [[PTR0_BEGIN:%.+]] = bitcast i8** [[PTR0]] to i32**
153 // CHECK: store i32* [[A_ADDR]], i32** [[PTR0_BEGIN]],
154 // CHECK: [[SIZE0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0
155 // CHECK: store i64 [[SZ]], i64* [[SIZE0]],
156 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 1
157 // CHECK: [[BPTR1_A_ADDR:%.+]] = bitcast i8** [[BPTR1]] to i32**
158 // CHECK: store i32* [[A_ADDR2]], i32** [[BPTR1_A_ADDR]],
159 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 1
160 // CHECK: [[PTR1_A_ADDR:%.+]] = bitcast i8** [[PTR1]] to i32**
161 // CHECK: store i32* [[A_ADDR2]], i32** [[PTR1_A_ADDR]],
162 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 2
163 // CHECK: [[BPTR2_PTR_ADDR:%.+]] = bitcast i8** [[BPTR2]] to i32***
164 // CHECK: store i32** [[PTR_ADDR]], i32*** [[BPTR2_PTR_ADDR]],
165 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 2
166 // CHECK: [[PTR2_PTR_ADDR:%.+]] = bitcast i8** [[PTR2]] to i32***
167 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR2_PTR_ADDR]],
168 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 3
169 // CHECK: [[BPTR3_REF_PTR:%.+]] = bitcast i8** [[BPTR3]] to i32**
170 // CHECK: store i32* [[REF_PTR]], i32** [[BPTR3_REF_PTR]],
171 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 3
172 // CHECK: [[PTR3_REF_PTR:%.+]] = bitcast i8** [[PTR3]] to i32**
173 // CHECK: store i32* [[REF_PTR]], i32** [[PTR3_REF_PTR]],
174 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 4
175 // CHECK: [[BPTR4_ARR_ADDR:%.+]] = bitcast i8** [[BPTR4]] to [4 x i32]**
176 // CHECK: store [4 x i32]* [[ARR_ADDR2]], [4 x i32]** [[BPTR4_ARR_ADDR]],
177 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 4
178 // CHECK: [[PTR4_ARR_ADDR:%.+]] = bitcast i8** [[PTR4]] to [4 x i32]**
179 // CHECK: store [4 x i32]* [[ARR_ADDR2]], [4 x i32]** [[PTR4_ARR_ADDR]],
180 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
181 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
182 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0
183 // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null)
184 // CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[BPTR1_A_ADDR]],
185 // CHECK: store i32* [[A_ADDR]], i32** [[A_REF:%.+]],
186 // CHECK: [[PTR_ADDR:%.+]] = load i32**, i32*** [[BPTR2_PTR_ADDR]],
187 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR_REF:%.+]],
188 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[BPTR3_REF_PTR]],
189 // CHECK: store i32* [[REF_PTR]], i32** [[REF_REF:%.+]],
190 // CHECK: [[PTR_ADDR:%.+]] = load i32**, i32*** [[BPTR2_PTR_ADDR]],
191 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR_REF2:%.+]],
192 // CHECK: [[ARR_ADDR:%.+]] = load [4 x i32]*, [4 x i32]** [[BPTR4_ARR_ADDR]],
193 // CHECK: store [4 x i32]* [[ARR_ADDR]], [4 x i32]** [[ARR_REF:%.+]],
194 // CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[A_REF]],
195 // CHECK: [[A:%.+]] = load i32, i32* [[A_ADDR]],
196 // CHECK: [[INC:%.+]] = add nsw i32 [[A]], 1
197 // CHECK: store i32 [[INC]], i32* [[A_ADDR]],
198 // CHECK: [[PTR_PTR:%.+]] = load i32**, i32*** [[PTR_REF2]],
199 // CHECK: [[PTR:%.+]] = load i32*, i32** [[PTR_PTR]],
200 // CHECK: [[VAL:%.+]] = load i32, i32* [[PTR]],
201 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
202 // CHECK: store i32 [[INC]], i32* [[PTR]],
203 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[REF_REF]],
204 // CHECK: [[VAL:%.+]] = load i32, i32* [[REF_PTR]],
205 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
206 // CHECK: store i32 [[INC]], i32* [[REF_PTR]],
207 // CHECK: [[ARR_ADDR:%.+]] = load [4 x i32]*, [4 x i32]** [[ARR_REF]],
208 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x i32], [4 x i32]* [[ARR_ADDR]], i64 0, i64 0
209 // CHECK: [[VAL:%.+]] = load i32, i32* [[ARR0_ADDR]],
210 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
211 // CHECK: store i32 [[INC]], i32* [[ARR0_ADDR]],
212 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
213 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
214 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0
215 // CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null)
216
217 #endif
218