Skip to content

Commit 0f8e7b4

Browse files
animeshk-amdanimeshk-amd
committed
[OpenMP] Add map clause to the LIT test on use_device_addr clause
As per the OpenMP Spec, "A list item in a use_device_addr clause must have a corresponding list item in the device data environment" . Therefore a `map` clause is added which will make sure that the respective list items are mapped to the device data environment before the `use_device_addr` clause is specified. The CHECK lines are also modified based on this change. Differential Revision: https://reviews.llvm.org/D134974
1 parent 4db6871 commit 0f8e7b4

File tree

1 file changed

+107
-83
lines changed

1 file changed

+107
-83
lines changed

‎clang/test/OpenMP/target_data_use_device_addr_codegen.cpp

Lines changed: 107 additions & 83 deletions
Original file line numberDiff line numberDiff line change
@@ -11,21 +11,21 @@
1111
#ifndef HEADER
1212
#define HEADER
1313

14-
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
14+
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 4, i64 16, i64 4, i64 4, i64 0, i64 4]
1515
// 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
16+
// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 67, i64 67, i64 3, i64 67, i64 67, i64 67]
17+
// CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 4, i64 16, i64 4, i64 4, i64 0]
1818
// 0 = OMP_MAP_NONE
1919
// 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]
20+
// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 281474976710723, i64 281474976710739, i64 281474976710739, i64 281474976710675, i64 281474976710723]
2121
struct S {
2222
int a = 0;
2323
int *ptr = &a;
2424
int &ref = a;
2525
int arr[4];
2626
S() {}
2727
void foo() {
28-
#pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a])
28+
#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:a]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a])
2929
++a, ++*ptr, ++ref, ++arr[0];
3030
}
3131
};
@@ -38,7 +38,7 @@ int main() {
3838
float vla[(int)a];
3939
S s;
4040
s.foo();
41-
#pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0])
41+
#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0])
4242
++a, ++*ptr, ++ref, ++arr[0], ++vla[0];
4343
return a;
4444
}
@@ -48,47 +48,68 @@ int main() {
4848
// CHECK: [[PTR_ADDR:%.+]] = alloca ptr,
4949
// CHECK: [[REF_ADDR:%.+]] = alloca ptr,
5050
// CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float],
51-
// CHECK: [[BPTRS:%.+]] = alloca [5 x ptr],
52-
// CHECK: [[PTRS:%.+]] = alloca [5 x ptr],
51+
// CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
52+
// CHECK: [[PTRS:%.+]] = alloca [6 x ptr],
53+
// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr],
54+
// CHECK: [[SIZES:%.+]] = alloca [6 x i64],
5355
// CHECK: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}},
5456
// CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_ADDR]],
55-
// CHECK: [[REF:%.+]] = load ptr, ptr [[REF_ADDR]],
56-
// CHECK: [[ARR:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_ADDR]], i64 0, i64 0
57-
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
57+
// CHECK-NEXT: [[P4:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
58+
// CHECK-NEXT: [[ARR_IDX:%.+]] = getelementptr inbounds float, ptr [[P4]], i64 3
59+
// CHECK: [[P5:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
60+
// CHECK-NEXT: [[P6:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
61+
// CHECK-NEXT: [[ARR_IDX1:%.+]] = getelementptr inbounds float, ptr [[P6]], i64 0
62+
// CHECK: [[P7:%.+]] = load ptr, ptr [[REF_ADDR]],
63+
// CHECK-NEXT: [[REF:%.+]] = load ptr, ptr [[REF_ADDR]],
64+
// CHECK-NEXT: [[ARR_IDX2:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_ADDR]], i64 0, i64 0
65+
// CHECK: [[P10:%.+]] = mul nuw i64 {{.+}}, 4
66+
// CHECK-NEXT: [[ARR_IDX5:%.+]] = getelementptr inbounds float, ptr [[VLA_ADDR]], i64 0
67+
// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[SIZES]], ptr align 8 [[SIZES1]], i64 48, i1 false)
68+
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
5869
// CHECK: store ptr [[A_ADDR]], ptr [[BPTR0]],
59-
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0
70+
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
6071
// CHECK: store ptr [[A_ADDR]], ptr [[PTR0]],
61-
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 1
72+
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1
6273
// CHECK: store ptr [[PTR]], ptr [[BPTR1]],
63-
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 1
64-
// CHECK: store ptr [[PTR]], ptr [[PTR1]],
65-
// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 2
66-
// CHECK: store ptr [[REF]], ptr [[BPTR2]],
67-
// CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 2
68-
// CHECK: store ptr [[REF]], ptr [[PTR2]],
69-
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 3
70-
// CHECK: store ptr [[ARR]], ptr [[BPTR3]],
71-
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 3
72-
// CHECK: store ptr [[ARR]], ptr [[PTR3]],
73-
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 4
74-
// CHECK: store ptr [[VLA_ADDR]], ptr [[BPTR4]],
75-
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 4
76-
// CHECK: store ptr [[VLA_ADDR]], ptr [[PTR4]],
77-
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
78-
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0
79-
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZES1]], ptr [[MAPTYPES1]], ptr null, ptr null)
74+
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1
75+
// CHECK: store ptr [[ARR_IDX]], ptr [[PTR1]],
76+
// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2
77+
// CHECK: store ptr [[P5]], ptr [[BPTR2]],
78+
// CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2
79+
// CHECK: store ptr [[ARR_IDX1]], ptr [[PTR2]],
80+
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3
81+
// CHECK: store ptr [[P7]], ptr [[BPTR3]],
82+
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3
83+
// CHECK: store ptr [[REF]], ptr [[PTR3]],
84+
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4
85+
// CHECK: store ptr [[ARR_ADDR]], ptr [[BPTR4]], align
86+
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4
87+
// CHECK: store ptr [[ARR_IDX2]], ptr [[PTR4]], align 8
88+
// CHECK: [[SIZE_PTR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 4
89+
// CHECK: store i64 [[P10:%.+]], ptr [[SIZE_PTR]], align 8
90+
// CHECK: [[MAP_PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[MAP_PTRS]], i64 0, i64 4
91+
// CHECK: store ptr null, ptr [[MAP_PTR]], align 8
92+
// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5
93+
// CHECK: store ptr [[VLA_ADDR]], ptr [[BPTR5]],
94+
// CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 5
95+
// CHECK: store ptr [[ARR_IDX5]], ptr [[PTR5]],
96+
97+
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
98+
// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
99+
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
100+
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null)
80101
// CHECK: [[A_REF:%.+]] = load ptr, ptr [[BPTR0]],
81-
// CHECK: [[REF_REF:%.+]] = load ptr, ptr [[BPTR2]],
102+
// CHECK: [[REF_REF:%.+]] = load ptr, ptr [[BPTR3]],
82103
// CHECK: store ptr [[REF_REF]], ptr [[TMP_REF_ADDR:%.+]],
83-
// CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR3]],
84-
// CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR4]],
104+
// CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR4]],
105+
// CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR5]],
85106
// CHECK: [[A:%.+]] = load float, ptr [[A_REF]],
86107
// CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00
87108
// CHECK: store float [[INC]], ptr [[A_REF]],
88-
// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR1]],
89-
// CHECK: [[VAL:%.+]] = load float, ptr [[PTR_ADDR]],
109+
// CHECK: [[PTR:%.+]] = load ptr, ptr [[BPTR1]],
110+
// CHECK: [[VAL:%.+]] = load float, ptr [[PTR]],
90111
// CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00
91-
// CHECK: store float [[INC]], ptr [[PTR_ADDR]],
112+
// CHECK: store float [[INC]], ptr [[PTR]],
92113
// CHECK: [[REF_ADDR:%.+]] = load ptr, ptr [[TMP_REF_ADDR]],
93114
// CHECK: [[REF:%.+]] = load float, ptr [[REF_ADDR]],
94115
// CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00
@@ -101,63 +122,66 @@ int main() {
101122
// CHECK: [[VLA0:%.+]] = load float, ptr [[VLA0_ADDR]],
102123
// CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00
103124
// CHECK: store float [[INC]], ptr [[VLA0_ADDR]],
104-
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
105-
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0
106-
// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZES1]], ptr [[MAPTYPES1]], ptr null, ptr null)
125+
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
126+
// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
127+
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
128+
// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null)
107129

108130
// CHECK: foo
109-
// %this.addr = alloca ptr, align 8
110-
// CHECK: [[BPTRS:%.+]] = alloca [5 x ptr],
111-
// CHECK: [[PTRS:%.+]] = alloca [5 x ptr],
112-
// CHECK: [[SIZES:%.+]] = alloca [5 x i64],
113-
// %tmp = alloca ptr, align 8
114-
// %tmp6 = alloca ptr, align 8
115-
// %tmp7 = alloca ptr, align 8
116-
// %tmp8 = alloca ptr, align 8
117-
// %tmp9 = alloca ptr, align 8
118-
// store ptr %this, ptr %this.addr, align 8
119-
// %this1 = load ptr, ptr %this.addr, align 8
131+
// CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
132+
// CHECK: [[PTRS:%.+]] = alloca [6 x ptr],
133+
// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr],
134+
// CHECK: [[SIZES:%.+]] = alloca [6 x i64],
120135
// CHECK: [[A_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS:%.+]], i32 0, i32 0
121-
// %ptr = getelementptr inbounds %struct.S, ptr %this1, i32 0, i32 1
122-
// %ref = getelementptr inbounds %struct.S, ptr %this1, i32 0, i32 2
123-
// %0 = load ptr, ptr %ref, align 8
124-
// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3
125-
// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 0
126136
// CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 1
137+
// CHECK: [[ARR_IDX:%.+]] = getelementptr inbounds i32, ptr %{{.+}}, i64 3
127138
// CHECK: [[REF_REF:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 2
128139
// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]],
129-
// CHECK: [[ARR_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3
130-
// CHECK: [[ARR_END:%.+]] = getelementptr [4 x i32], ptr [[ARR_ADDR]], i32 1
140+
// CHECK-NEXT: [[P3:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 1
141+
// CHECK: [[ARR_IDX5:%.+]] = getelementptr inbounds i32, ptr {{.+}}, i64 0
142+
// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3
143+
144+
// CHECK: [[ARR_IDX6:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0
145+
// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 0
146+
// CHECK: [[P4:%.+]] = mul nuw i64 [[CONV:%.+]], 4
147+
// CHECK: [[ARR_END:%.+]] = getelementptr i32, ptr [[ARR_IDX6]], i32 1
131148
// CHECK: [[E:%.+]] = ptrtoint ptr [[ARR_END]] to i64
132149
// CHECK: [[B:%.+]] = ptrtoint ptr [[A_ADDR]] to i64
133150
// CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]]
134151
// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
135-
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
152+
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
136153
// CHECK: store ptr [[THIS]], ptr [[BPTR0]],
137-
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0
154+
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
138155
// CHECK: store ptr [[A_ADDR]], ptr [[PTR0]],
139-
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [5 x i64], ptr [[SIZES]], i32 0, i32 0
156+
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
140157
// CHECK: store i64 [[SZ]], ptr [[SIZE0]],
141-
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 1
142-
// CHECK: store ptr [[A_ADDR2]], ptr [[BPTR1]],
143-
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 1
144-
// CHECK: store ptr [[A_ADDR2]], ptr [[PTR1]],
145-
// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 2
158+
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1
159+
// CHECK: store ptr [[THIS]], ptr [[BPTR1]]
160+
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1
161+
// CHECK: store ptr [[A_ADDR]], ptr [[PTR1]],
162+
// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2
146163
// CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]],
147-
// CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 2
148-
// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR2]],
149-
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 3
150-
// CHECK: store ptr [[REF_PTR]], ptr [[BPTR3]],
151-
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 3
164+
// CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2
165+
// CHECK: store ptr [[ARR_IDX]], ptr [[PTR2]],
166+
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3
167+
// CHECK: store ptr [[THIS]], ptr [[BPTR3]]
168+
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3
152169
// CHECK: store ptr [[REF_PTR]], ptr [[PTR3]],
153-
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 4
154-
// CHECK: store ptr [[ARR_ADDR2]], ptr [[BPTR4]],
155-
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 4
156-
// CHECK: store ptr [[ARR_ADDR2]], ptr [[PTR4]],
157-
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
158-
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0
159-
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], ptr [[SIZES]], i32 0, i32 0
160-
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null)
170+
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4
171+
// CHECK: store ptr [[P3]], ptr [[BPTR4]],
172+
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4
173+
// CHECK: store ptr [[ARR_IDX5]], ptr [[PTR4]]
174+
175+
// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5
176+
// CHECK: store ptr [[THIS]], ptr [[BPTR5]], align 8
177+
// CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 5
178+
// CHECK: store ptr [[ARR_IDX6]], ptr [[PTR5]], align 8
179+
// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 5
180+
// CHECK: store i64 [[P4]], ptr [[SIZE1]], align 8
181+
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
182+
// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
183+
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
184+
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null)
161185
// CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[BPTR1]],
162186
// CHECK: store ptr [[A_ADDR]], ptr [[A_REF:%.+]],
163187
// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]],
@@ -166,7 +190,7 @@ int main() {
166190
// CHECK: store ptr [[REF_PTR]], ptr [[REF_REF:%.+]],
167191
// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]],
168192
// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF2:%.+]],
169-
// CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[BPTR4]],
193+
// CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[BPTR5]],
170194
// CHECK: store ptr [[ARR_ADDR]], ptr [[ARR_REF:%.+]],
171195
// CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[A_REF]],
172196
// CHECK: [[A:%.+]] = load i32, ptr [[A_ADDR]],
@@ -186,9 +210,9 @@ int main() {
186210
// CHECK: [[VAL:%.+]] = load i32, ptr [[ARR0_ADDR]],
187211
// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
188212
// CHECK: store i32 [[INC]], ptr [[ARR0_ADDR]],
189-
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
190-
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0
191-
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], ptr [[SIZES]], i32 0, i32 0
192-
// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null)
213+
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
214+
// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
215+
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
216+
// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null)
193217

194218
#endif

0 commit comments

Comments
 (0)