Skip to content

Commit

Permalink
[OPENMP]Fix PR46170: partial mapping for array sections of data members.
Browse files Browse the repository at this point in the history 
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
  • Loading branch information
@alexey-bataev
alexey-bataev committed Jun 3, 2020
1 parent c1911fc commit 59e0987
Show file tree
Hide file tree
Showing 2 changed files with 122 additions and 1 deletion.
9 changes: 8 additions & 1 deletion clang/lib/CodeGen/CGOpenMPRuntime.cpp
View file
Expand Up @@ -7736,7 +7736,14 @@ class MappableExprsHandler {
// Update info about the lowest and highest elements for this struct
if (!PartialStruct.Base.isValid()) {
PartialStruct.LowestElem = {FieldIndex, LB};
PartialStruct.HighestElem = {FieldIndex, LB};
if (IsFinalArraySection) {
Address HB =
CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
.getAddress(CGF);
PartialStruct.HighestElem = {FieldIndex, HB};
} else {
PartialStruct.HighestElem = {FieldIndex, LB};
}
PartialStruct.Base = BP;
} else if (FieldIndex < PartialStruct.LowestElem.first) {
PartialStruct.LowestElem = {FieldIndex, LB};
Expand Down
114 changes: 114 additions & 0 deletions clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
View file
@@ -0,0 +1,114 @@
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// 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
// expected-no-diagnostics
#ifndef HEADER
#define HEADER

// 32 = 0x20 = OMP_MAP_TARGET_PARAM
// 281474976710656 = 0x1000000000000 = OMP_MAP_MEMBER_OF of 1-st element
// CHECK: [[MAP_ENTER:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710656]
// 281474976710664 = 0x1000000000008 = OMP_MAP_MEMBER_OF of 1-st element | OMP_MAP_DELETE
// CHECK: [[MAP_EXIT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710664]
template <typename T>
struct S {
constexpr static int size = 6;
T data[size];
};

template <typename T>
struct maptest {
S<T> s;
maptest() {
// CHECK: [[BPTRS:%.+]] = alloca [2 x i8*],
// CHECK: [[PTRS:%.+]] = alloca [2 x i8*],
// CHECK: [[SIZES:%.+]] = alloca [2 x i64],
// CHECK: getelementptr inbounds
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS:%.+]], i32 0, i32 0
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
// CHECK: [[S_DATA_0_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 0

// SZ = &this->s.data[6]-&this->s.data[0]
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS]], i32 0, i32 0
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
// CHECK: [[S_DATA_5_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 5
// CHECK: [[S_DATA_6_ADDR:%.+]] = getelementptr float, float* [[S_DATA_5_ADDR]], i32 1
// CHECK: [[BEG:%.+]] = bitcast float* [[S_DATA_0_ADDR]] to i8*
// CHECK: [[END:%.+]] = bitcast float* [[S_DATA_6_ADDR]] to i8*
// CHECK: [[END_BC:%.+]] = ptrtoint i8* [[END]] to i64
// CHECK: [[BEG_BC:%.+]] = ptrtoint i8* [[BEG]] to i64
// CHECK: [[DIFF:%.+]] = sub i64 [[END_BC]], [[BEG_BC]]
// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)

// Fill mapping arrays
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
// CHECK: [[BPTR0_THIS:%.+]] = bitcast i8** [[BPTR0]] to %struct.maptest**
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR0_THIS]],
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
// CHECK: [[PTR0_DATA:%.+]] = bitcast i8** [[PTR0]] to float**
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR0_DATA]],
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
// CHECK: store i64 [[SZ]], i64* [[SIZE0]],
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 1
// CHECK: [[BPTR1_THIS:%.+]] = bitcast i8** [[BPTR1]] to %struct.maptest**
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR1_THIS]],
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 1
// CHECK: [[PTR1_DATA:%.+]] = bitcast i8** [[PTR1]] to float**
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR1_DATA]],
// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 1
// CHECK: store i64 24, i64* [[SIZE1]],
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
// 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))
#pragma omp target enter data map(alloc : s.data[:6])
}

~maptest() {
// CHECK: [[BPTRS:%.+]] = alloca [2 x i8*],
// CHECK: [[PTRS:%.+]] = alloca [2 x i8*],
// CHECK: [[SIZE:%.+]] = alloca [2 x i64],
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS:%.+]], i32 0, i32 0
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
// CHECK: [[S_DATA_0_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 0

// SZ = &this->s.data[6]-&this->s.data[0]
// CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS]], i32 0, i32 0
// CHECK: [[S_DATA_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_ADDR]], i32 0, i32 0
// CHECK: [[S_DATA_5_ADDR:%.+]] = getelementptr inbounds [6 x float], [6 x float]* [[S_DATA_ADDR]], i64 0, i64 5
// CHECK: [[S_DATA_6_ADDR:%.+]] = getelementptr float, float* [[S_DATA_5_ADDR]], i32 1
// CHECK: [[BEG:%.+]] = bitcast float* [[S_DATA_0_ADDR]] to i8*
// CHECK: [[END:%.+]] = bitcast float* [[S_DATA_6_ADDR]] to i8*
// CHECK: [[END_BC:%.+]] = ptrtoint i8* [[END]] to i64
// CHECK: [[BEG_BC:%.+]] = ptrtoint i8* [[BEG]] to i64
// CHECK: [[DIFF:%.+]] = sub i64 [[END_BC]], [[BEG_BC]]
// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)

// Fill mapping arrays
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
// CHECK: [[BPTR0_THIS:%.+]] = bitcast i8** [[BPTR0]] to %struct.maptest**
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR0_THIS]],
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
// CHECK: [[PTR0_DATA:%.+]] = bitcast i8** [[PTR0]] to float**
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR0_DATA]],
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
// CHECK: store i64 [[SZ]], i64* [[SIZE0]],
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 1
// CHECK: [[BPTR1_THIS:%.+]] = bitcast i8** [[BPTR1]] to %struct.maptest**
// CHECK: store %struct.maptest* [[THIS]], %struct.maptest** [[BPTR1_THIS]],
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 1
// CHECK: [[PTR1_DATA:%.+]] = bitcast i8** [[PTR1]] to float**
// CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR1_DATA]],
// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 1
// CHECK: store i64 24, i64* [[SIZE1]],
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
// 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))
#pragma omp target exit data map(delete : s.data[:6])
}
};

maptest<float> a;

#endif

0 comments on commit 59e0987

Please sign in to comment.