[OpenMP][Map][NFC] improve map chain. #101903

jyu2-git · 2024-08-04T19:17:26Z

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.

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.

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.

llvmbot · 2024-08-05T05:08:51Z

@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: None (jyu2-git)

Changes

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.

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.

Full diff: https://github.com/llvm/llvm-project/pull/101903.diff

2 Files Affected:

(modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+14-8)
(added) clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp (+50)

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]

alexey-bataev

LG

jyu2-git · 2024-08-05T14:58:37Z

LG

Thanks, Alexey!!

jyu2-git force-pushed the partialmaptype branch from 6ed9ae5 to 318c3b4 Compare August 5, 2024 05:05

jyu2-git changed the title ~~Partialmaptype~~ [OpenMP][Map][NFC] improve map chain. Aug 5, 2024

jyu2-git marked this pull request as ready for review August 5, 2024 05:08

llvmbot added clang Clang issues not falling into any other category clang:codegen IR generation bugs: mangling, exceptions, etc. clang:openmp OpenMP related changes to Clang labels Aug 5, 2024

jyu2-git requested a review from alexey-bataev August 5, 2024 05:08

alexey-bataev approved these changes Aug 5, 2024

View reviewed changes

jyu2-git merged commit 6848b99 into llvm:main Aug 5, 2024
12 checks passed

jyu2-git added a commit to jyu2-git/llvm-project that referenced this pull request Aug 6, 2024

Merge branch 'llvm:main' into llvm#101903

224c047

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[OpenMP][Map][NFC] improve map chain. #101903

[OpenMP][Map][NFC] improve map chain. #101903

Uh oh!

jyu2-git commented Aug 4, 2024 •

edited

Loading

Uh oh!

llvmbot commented Aug 5, 2024 •

edited

Loading

Uh oh!

alexey-bataev left a comment

Uh oh!

jyu2-git commented Aug 5, 2024

Uh oh!

Uh oh!

Uh oh!

[OpenMP][Map][NFC] improve map chain. #101903

[OpenMP][Map][NFC] improve map chain. #101903

Uh oh!

Conversation

jyu2-git commented Aug 4, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Aug 5, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

alexey-bataev left a comment

Choose a reason for hiding this comment

Uh oh!

jyu2-git commented Aug 5, 2024

Uh oh!

Uh oh!

Uh oh!

jyu2-git commented Aug 4, 2024 •

edited

Loading

llvmbot commented Aug 5, 2024 •

edited

Loading