Skip to content

Commit 59e0987

Browse files
alexey-bataevalexey-bataev
committed
[OPENMP]Fix PR46170: partial mapping for array sections of data members.
Summary: If the data member is mapped as an array section, need to emit the pointer to the last element of this array section and use this pointer as the highest element in partial struct data. Reviewers: jdoerfert Subscribers: yaxunl, guansong, sstefan1, cfe-commits, caomhin Tags: #clang Differential Revision: https://reviews.llvm.org/D81037
1 parent c1911fc commit 59e0987

File tree

2 files changed

+122
-1
lines changed

2 files changed

+122
-1
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7736,7 +7736,14 @@ class MappableExprsHandler {
77367736
// Update info about the lowest and highest elements for this struct
77377737
if (!PartialStruct.Base.isValid()) {
77387738
PartialStruct.LowestElem = {FieldIndex, LB};
7739-
PartialStruct.HighestElem = {FieldIndex, LB};
7739+
if (IsFinalArraySection) {
7740+
Address HB =
7741+
CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
7742+
.getAddress(CGF);
7743+
PartialStruct.HighestElem = {FieldIndex, HB};
7744+
} else {
7745+
PartialStruct.HighestElem = {FieldIndex, LB};
7746+
}
77407747
PartialStruct.Base = BP;
77417748
} else if (FieldIndex < PartialStruct.LowestElem.first) {
77427749
PartialStruct.LowestElem = {FieldIndex, LB};

‎clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp

Lines changed: 114 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,114 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
2+
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
3+
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
4+
// expected-no-diagnostics
5+
#ifndef HEADER
6+
#define HEADER
7+
8+
// 32 = 0x20 = OMP_MAP_TARGET_PARAM
9+
// 281474976710656 = 0x1000000000000 = OMP_MAP_MEMBER_OF of 1-st element
10+
// CHECK: [[MAP_ENTER:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710656]
11+
// 281474976710664 = 0x1000000000008 = OMP_MAP_MEMBER_OF of 1-st element | OMP_MAP_DELETE
12+
// CHECK: [[MAP_EXIT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710664]
13+
template <typename T>
14+
struct S {
15+
constexpr static int size = 6;
16+
T data[size];
17+
};
18+
19+
template <typename T>
20+
struct maptest {
21+
S<T> s;
22+
maptest() {
23+
// CHECK: [[BPTRS:%.+]] = alloca [2 x i8*],
24+
// CHECK: [[PTRS:%.+]] = alloca [2 x i8*],
25+
// CHECK: [[SIZES:%.+]] = alloca [2 x i64],
26+
// CHECK: getelementptr inbounds
27+
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS:%.+]], i32 0, i32 0
28+
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
29+
// CHECK: [[S_DATA_0_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 0
30+
31+
// SZ = &this->s.data[6]-&this->s.data[0]
32+
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS]], i32 0, i32 0
33+
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
34+
// CHECK: [[S_DATA_5_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 5
35+
// CHECK: [[S_DATA_6_ADDR:%.+]] = getelementptr float, float* [[S_DATA_5_ADDR]], i32 1
36+
// CHECK: [[BEG:%.+]] = bitcast float* [[S_DATA_0_ADDR]] to i8*
37+
// CHECK: [[END:%.+]] = bitcast float* [[S_DATA_6_ADDR]] to i8*
38+
// CHECK: [[END_BC:%.+]] = ptrtoint i8* [[END]] to i64
39+
// CHECK: [[BEG_BC:%.+]] = ptrtoint i8* [[BEG]] to i64
40+
// CHECK: [[DIFF:%.+]] = sub i64 [[END_BC]], [[BEG_BC]]
41+
// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
42+
43+
// Fill mapping arrays
44+
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
45+
// CHECK: [[BPTR0_THIS:%.+]] = bitcast i8** [[BPTR0]] to %struct.maptest**
46+
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR0_THIS]],
47+
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
48+
// CHECK: [[PTR0_DATA:%.+]] = bitcast i8** [[PTR0]] to float**
49+
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR0_DATA]],
50+
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
51+
// CHECK: store i64 [[SZ]], i64* [[SIZE0]],
52+
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 1
53+
// CHECK: [[BPTR1_THIS:%.+]] = bitcast i8** [[BPTR1]] to %struct.maptest**
54+
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR1_THIS]],
55+
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 1
56+
// CHECK: [[PTR1_DATA:%.+]] = bitcast i8** [[PTR1]] to float**
57+
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR1_DATA]],
58+
// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 1
59+
// CHECK: store i64 24, i64* [[SIZE1]],
60+
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
61+
// CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
62+
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
63+
// CHECK: call void @__tgt_target_data_begin(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_ENTER]], i32 0, i32 0))
64+
#pragma omp target enter data map(alloc : s.data[:6])
65+
}
66+
67+
~maptest() {
68+
// CHECK: [[BPTRS:%.+]] = alloca [2 x i8*],
69+
// CHECK: [[PTRS:%.+]] = alloca [2 x i8*],
70+
// CHECK: [[SIZE:%.+]] = alloca [2 x i64],
71+
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS:%.+]], i32 0, i32 0
72+
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
73+
// CHECK: [[S_DATA_0_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 0
74+
75+
// SZ = &this->s.data[6]-&this->s.data[0]
76+
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS]], i32 0, i32 0
77+
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
78+
// CHECK: [[S_DATA_5_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 5
79+
// CHECK: [[S_DATA_6_ADDR:%.+]] = getelementptr float, float* [[S_DATA_5_ADDR]], i32 1
80+
// CHECK: [[BEG:%.+]] = bitcast float* [[S_DATA_0_ADDR]] to i8*
81+
// CHECK: [[END:%.+]] = bitcast float* [[S_DATA_6_ADDR]] to i8*
82+
// CHECK: [[END_BC:%.+]] = ptrtoint i8* [[END]] to i64
83+
// CHECK: [[BEG_BC:%.+]] = ptrtoint i8* [[BEG]] to i64
84+
// CHECK: [[DIFF:%.+]] = sub i64 [[END_BC]], [[BEG_BC]]
85+
// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
86+
87+
// Fill mapping arrays
88+
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
89+
// CHECK: [[BPTR0_THIS:%.+]] = bitcast i8** [[BPTR0]] to %struct.maptest**
90+
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR0_THIS]],
91+
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
92+
// CHECK: [[PTR0_DATA:%.+]] = bitcast i8** [[PTR0]] to float**
93+
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR0_DATA]],
94+
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
95+
// CHECK: store i64 [[SZ]], i64* [[SIZE0]],
96+
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 1
97+
// CHECK: [[BPTR1_THIS:%.+]] = bitcast i8** [[BPTR1]] to %struct.maptest**
98+
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR1_THIS]],
99+
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 1
100+
// CHECK: [[PTR1_DATA:%.+]] = bitcast i8** [[PTR1]] to float**
101+
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR1_DATA]],
102+
// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 1
103+
// CHECK: store i64 24, i64* [[SIZE1]],
104+
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
105+
// CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
106+
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
107+
// CHECK: call void @__tgt_target_data_end(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_EXIT]], i32 0, i32 0))
108+
#pragma omp target exit data map(delete : s.data[:6])
109+
}
110+
};
111+
112+
maptest<float> a;
113+
114+
#endif

0 commit comments

Comments
 (0)