https://github.com/jyu2-git updated
https://github.com/llvm/llvm-project/pull/101903
>From 318c3b401627acc158d320d3d5bb7c237d2e9da6 Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer...@intel.com>
Date: Sun, 4 Aug 2024 10:53:29 -0700
Subject: [PATCH] [OpenMP][Map][NFC] improve map chain.
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 intead 8.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 22 +++++---
...get_map_pointer_defalut_mapper_codegen.cpp | 50 +++++++++++++++++++
2 files changed, 64 insertions(+), 8 deletions(-)
create mode 100644
clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index d869aa3322cce..de04b7997ca37 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7125,6 +7125,9 @@ class MappableExprsHandler {
bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
bool IsPrevMemberReference = false;
+ bool IsPartialMapped =
+ !PartialStruct.PreliminaryMapData.BasePointers.empty();
+
// We need to check if we will be encountering any MEs. If we do not
// encounter any ME expression it means we will be mapping the whole
struct.
// In that case we need to skip adding an entry for the struct to the
@@ -7370,7 +7373,9 @@ class MappableExprsHandler {
// whole struct is currently being mapped. The struct needs to be added
// in the first position before any data internal to the struct is
being
// mapped.
- if (!IsMemberPointerOrAddr ||
+ // Skip adding an entry in the CurInfo of this combined entry if the
+ // PartialStruct.PreliminaryMapData.BasePointers has been mapped.
+ if ((!IsMemberPointerOrAddr && !IsPartialMapped) ||
(Next == CE && MapType != OMPC_MAP_unknown)) {
if (!IsMappingWholeStruct) {
CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
@@ -7486,8 +7491,8 @@ class MappableExprsHandler {
// The pointer becomes the base for the next element.
if (Next != CE)
BP = IsMemberReference ? LowestElem : LB;
-
- IsExpressionFirstInfo = false;
+ if (!IsPartialMapped)
+ IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
FirstPointerInComplexData = false;
IsPrevMemberReference = IsMemberReference;
@@ -8295,7 +8300,9 @@ class MappableExprsHandler {
// Map type is always TARGET_PARAM, if generate info for captures.
CombinedInfo.Types.push_back(
NotTargetParams ? OpenMPOffloadMappingFlags::OMP_MAP_NONE
- : OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
+ : !PartialStruct.PreliminaryMapData.BasePointers.empty()
+ ? OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ
+ : OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
// If any element has the present modifier, then make sure the runtime
// doesn't attempt to allocate the struct.
if (CurTypes.end() !=
@@ -9525,10 +9532,9 @@ static void genMapInfoForCaptures(
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid()) {
CombinedInfo.append(PartialStruct.PreliminaryMapData);
- MEHandler.emitCombinedEntry(
- CombinedInfo, CurInfo.Types, PartialStruct, CI->capturesThis(),
- OMPBuilder, nullptr,
- !PartialStruct.PreliminaryMapData.BasePointers.empty());
+ MEHandler.emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
+ CI->capturesThis(), OMPBuilder, nullptr,
+ /*NotTargetParams*/ false);
}
// We need to append the results of this capture to what we already have.
diff --git a/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
b/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
new file mode 100644
index 0000000000000..6ab10c45beb25
--- /dev/null
+++ b/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
@@ -0,0 +1,50 @@
+// 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
+
+typedef struct {
+ int a;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a)
+
+typedef struct {
+ int e;
+ C f;
+ int h;
+} D;
+
+D s;
+
+void foo() {
+ s.e = 111;
+ s.h = 10;
+ D *sp = &s;
+ D **spp = &sp;
+ D ***sppp = &spp;
+#pragma omp target map(tofrom : spp[0][0])
+ {
+ spp[0][0].e = 333;
+ }
+#pragma omp target map(tofrom : sp[0])
+ {
+ sp[0].e = 444;
+ }
+#pragma omp target map(tofrom : sppp[0][0][0])
+ {
+ sppp[0][0][0].e = 555;
+ }
+}
+#endif
+
+// CHECK: @.offload_sizes = private unnamed_addr constant [5 x i64] [i64 8,
i64 0, i64 0, i64 0, i64 4]
+// CHECK-NOT: @.offload_sizes = private unnamed_addr constant [6 x i64] [i64
8, i64 0, i64 0, i64 0, i64 8, i64 4]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [5 x i64] [i64
35, i64 16, i64 562949953421315, i64 562949953421315, i64 562949953421827]
+// CHECK-NOT: .offload_maptypes = private unnamed_addr constant [6 x i64]
[i64 35, i64 0, i64 562949953421315, i64 562949953421315, i64 562949953421827,
i64 562949953421843]
+// CHECK: @.offload_sizes.1 = private unnamed_addr constant [4 x i64] [i64 0,
i64 0, i64 0, i64 4]
+// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [4 x i64] [i64
32, i64 281474976710659, i64 281474976710659, i64 281474976711171]
+// CHECK: @.offload_sizes.3 = private unnamed_addr constant [6 x i64] [i64 8,
i64 8, i64 0, i64 0, i64 0, i64 4]
+// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [6 x i64] [i64
35, i64 16, i64 16, i64 844424930131971, i64 844424930131971, i64
844424930132483]
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
