Skip to content

Commit 6848b99

Browse files
authored
[OpenMP][Map][NFC] improve map chain. (llvm#101903)
This is for mapping structure has data members, which have 'default' mappers, where needs to map these members individually using their 'default' mappers. example map(tofrom: spp[0][0]), look at test case. currently create 6 maps: 1>&spp, &spp[0], size 8, maptype TARGET_PARAM | FROM | TO 2>&spp[0], &spp[0][0], size(D)with maptype OMP_MAP_NONE, nullptr 3>&spp[0], &spp[0][0].e, size(e) with maptype MEMBER_OF | FROM | TO 4>&spp[0], &spp[0][0].h, size(h) with maptype MEMBER_OF | FROM | TO 5>&spp, &spp[0],size(8), maptype MEMBER_OF | IMPLICIT | FROM | TO 6>&spp[0], &spp[0][0].f size(D) with maptype MEMBER_OF |IMPLICIT |PTR_AND_OBJ, @.omp_mapper._ZTS1C.default maptype with/without OMP_MAP_PTR_AND_OBJ For "2" and "5", since it is mapping pointer and pointee pair, PTR_AND_OBJ should be set But for "6" the PTR_AND_OBJ should not set. However, "5" is duplicate with "1" can be skip. To fix "2", during the call to emitCombinEntry with false with NotTargetParams instead !PartialStruct.PreliminaryMapData.BasePointers.empty(), since all captures need to be TARGET_PARAM And inside emitCombineEntry: check !PartialStruct.PreliminaryMapData.BasePointers.empty() to set PTR_AND_OBJ For "5" and "6": the fix in generateInfoForComponentList: Add new variable IsPartialMapped set with !PartialStruct.PreliminaryMapData.BasePointers.empty(); When that is true, skip generate "5" and don"t set IsExpressionFirstInfo to false, so that PTR_AND_OBJ would be set. After fix: will have 5 maps instead 6 1>&spp, &spp[0], size 8, maptype TARGET_PARAM | FROM | TO 2>&spp[0], &spp[0][0], size(D), maptype PTR_AND_OBJ, nullptr 3>&spp[0], &spp[0][0].e, size(e), maptype MEMBER_OF_2 | FROM | TO 4>&spp[0], &spp[0][0].h, size(h), maptype MEMBER_OF_2 | FROM | TO 5>&spp[0], &spp[0][0].f size(32), maptype MEMBER_OF_2 | IMPLICIT, @.omp_mapper._ZTS1C.default For map(sppp[0][0][0]): after fix: will have 6 maps instead 8. llvm#101903
1 parent 337c8b1 commit 6848b99

File tree

2 files changed

+64
-8
lines changed

2 files changed

+64
-8
lines changed

clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 14 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -7125,6 +7125,9 @@ class MappableExprsHandler {
71257125
bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
71267126
bool IsPrevMemberReference = false;
71277127

7128+
bool IsPartialMapped =
7129+
!PartialStruct.PreliminaryMapData.BasePointers.empty();
7130+
71287131
// We need to check if we will be encountering any MEs. If we do not
71297132
// encounter any ME expression it means we will be mapping the whole struct.
71307133
// In that case we need to skip adding an entry for the struct to the
@@ -7370,7 +7373,9 @@ class MappableExprsHandler {
73707373
// whole struct is currently being mapped. The struct needs to be added
73717374
// in the first position before any data internal to the struct is being
73727375
// mapped.
7373-
if (!IsMemberPointerOrAddr ||
7376+
// Skip adding an entry in the CurInfo of this combined entry if the
7377+
// PartialStruct.PreliminaryMapData.BasePointers has been mapped.
7378+
if ((!IsMemberPointerOrAddr && !IsPartialMapped) ||
73747379
(Next == CE && MapType != OMPC_MAP_unknown)) {
73757380
if (!IsMappingWholeStruct) {
73767381
CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
@@ -7486,8 +7491,8 @@ class MappableExprsHandler {
74867491
// The pointer becomes the base for the next element.
74877492
if (Next != CE)
74887493
BP = IsMemberReference ? LowestElem : LB;
7489-
7490-
IsExpressionFirstInfo = false;
7494+
if (!IsPartialMapped)
7495+
IsExpressionFirstInfo = false;
74917496
IsCaptureFirstInfo = false;
74927497
FirstPointerInComplexData = false;
74937498
IsPrevMemberReference = IsMemberReference;
@@ -8295,7 +8300,9 @@ class MappableExprsHandler {
82958300
// Map type is always TARGET_PARAM, if generate info for captures.
82968301
CombinedInfo.Types.push_back(
82978302
NotTargetParams ? OpenMPOffloadMappingFlags::OMP_MAP_NONE
8298-
: OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
8303+
: !PartialStruct.PreliminaryMapData.BasePointers.empty()
8304+
? OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ
8305+
: OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
82998306
// If any element has the present modifier, then make sure the runtime
83008307
// doesn't attempt to allocate the struct.
83018308
if (CurTypes.end() !=
@@ -9525,10 +9532,9 @@ static void genMapInfoForCaptures(
95259532
// individual members mapped. Emit an extra combined entry.
95269533
if (PartialStruct.Base.isValid()) {
95279534
CombinedInfo.append(PartialStruct.PreliminaryMapData);
9528-
MEHandler.emitCombinedEntry(
9529-
CombinedInfo, CurInfo.Types, PartialStruct, CI->capturesThis(),
9530-
OMPBuilder, nullptr,
9531-
!PartialStruct.PreliminaryMapData.BasePointers.empty());
9535+
MEHandler.emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
9536+
CI->capturesThis(), OMPBuilder, nullptr,
9537+
/*NotTargetParams*/ false);
95329538
}
95339539

95349540
// We need to append the results of this capture to what we already have.
Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,50 @@
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+
5+
// expected-no-diagnostics
6+
#ifndef HEADER
7+
#define HEADER
8+
9+
typedef struct {
10+
int a;
11+
} C;
12+
#pragma omp declare mapper(C s) map(to : s.a)
13+
14+
typedef struct {
15+
int e;
16+
C f;
17+
int h;
18+
} D;
19+
20+
D s;
21+
22+
void foo() {
23+
s.e = 111;
24+
s.h = 10;
25+
D *sp = &s;
26+
D **spp = &sp;
27+
D ***sppp = &spp;
28+
#pragma omp target map(tofrom : spp[0][0])
29+
{
30+
spp[0][0].e = 333;
31+
}
32+
#pragma omp target map(tofrom : sp[0])
33+
{
34+
sp[0].e = 444;
35+
}
36+
#pragma omp target map(tofrom : sppp[0][0][0])
37+
{
38+
sppp[0][0][0].e = 555;
39+
}
40+
}
41+
#endif
42+
43+
// CHECK: @.offload_sizes = private unnamed_addr constant [5 x i64] [i64 8, i64 0, i64 0, i64 0, i64 4]
44+
// CHECK-NOT: @.offload_sizes = private unnamed_addr constant [6 x i64] [i64 8, i64 0, i64 0, i64 0, i64 8, i64 4]
45+
// CHECK: @.offload_maptypes = private unnamed_addr constant [5 x i64] [i64 35, i64 16, i64 562949953421315, i64 562949953421315, i64 562949953421827]
46+
// CHECK-NOT: .offload_maptypes = private unnamed_addr constant [6 x i64] [i64 35, i64 0, i64 562949953421315, i64 562949953421315, i64 562949953421827, i64 562949953421843]
47+
// CHECK: @.offload_sizes.1 = private unnamed_addr constant [4 x i64] [i64 0, i64 0, i64 0, i64 4]
48+
// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [4 x i64] [i64 32, i64 281474976710659, i64 281474976710659, i64 281474976711171]
49+
// CHECK: @.offload_sizes.3 = private unnamed_addr constant [6 x i64] [i64 8, i64 8, i64 0, i64 0, i64 0, i64 4]
50+
// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [6 x i64] [i64 35, i64 16, i64 16, i64 844424930131971, i64 844424930131971, i64 844424930132483]

0 commit comments

Comments
 (0)