[Openmp-commits] [openmp] [OpenMP][libomptarget][FIX] Fix mapping of struct to device (PR #70821)

Gheorghe-Teodor Bercea via Openmp-commits openmp-commits at lists.llvm.org
Tue Nov 14 11:47:34 PST 2023


https://github.com/doru1004 updated https://github.com/llvm/llvm-project/pull/70821

>From e8d75c4dd9243894340504ee61609507b8e552c5 Mon Sep 17 00:00:00 2001
From: Doru Bercea <doru.bercea at amd.com>
Date: Tue, 31 Oct 2023 11:59:36 -0400
Subject: [PATCH] Fix mapping of struct.

---
 openmp/libomptarget/src/omptarget.cpp         |  61 +++++++---
 .../struct_mapping_with_pointers.cpp          | 114 ++++++++++++++++++
 2 files changed, 158 insertions(+), 17 deletions(-)
 create mode 100644 openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp

diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 6c59bc1cf38a8bb..eb9e5abe00b4609 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -585,11 +585,38 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
                     void **ArgMappers, AsyncInfoTy &AsyncInfo,
                     bool FromMapper) {
   TIMESCOPE_WITH_IDENT(Loc);
+
+  // Initialize new map type with old type:
+  SmallVector<int64_t, 16> NewArgTypes(ArgTypes, ArgTypes + ArgNum);
+
+  // Try to prevent mapping a struct multiple times in the same construct.
+  // Mapping the struct more than once will potentially overwrite previously
+  // mapped information.
+  for (int32_t I = 0; I < ArgNum; ++I) {
+    if (NewArgTypes[I] < 0)
+      continue;
+    if ((NewArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
+        (NewArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
+      continue;
+    for (int32_t J = I + 1; J < ArgNum; ++J) {
+      if (Args[I] == ArgsBase[I] && Args[I] == Args[J] &&
+          ArgsBase[I] == ArgsBase[J] && ArgSizes[I] == ArgSizes[J] &&
+          ArgSizes[I] > 0 && NewArgTypes[J] >= 0) {
+        NewArgTypes[I] |= ArgTypes[J];
+        NewArgTypes[J] = -1;
+      }
+    }
+  }
+
   // process each input.
   for (int32_t I = 0; I < ArgNum; ++I) {
+    int64_t ArgType = NewArgTypes[I];
+    if (ArgType < 0)
+      continue;
+
     // Ignore private variables and arrays - there is no mapping for them.
-    if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
-        (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
+    if ((ArgType & OMP_TGT_MAPTYPE_LITERAL) ||
+        (ArgType & OMP_TGT_MAPTYPE_PRIVATE))
       continue;
 
     if (ArgMappers && ArgMappers[I]) {
@@ -600,7 +627,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
 
       map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
       int Rc = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],
-                                ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,
+                                ArgType, ArgName, ArgMappers[I], AsyncInfo,
                                 targetDataBegin);
 
       if (Rc != OFFLOAD_SUCCESS) {
@@ -623,8 +650,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     // is a combined entry.
     int64_t TgtPadding = 0;
     const int NextI = I + 1;
-    if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
-        getParentIndex(ArgTypes[NextI]) == I) {
+    if (getParentIndex(ArgType) < 0 && NextI < ArgNum &&
+        NewArgTypes[NextI] >= 0 && getParentIndex(NewArgTypes[NextI]) == I) {
       int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);
       TgtPadding = (int64_t)HstPtrBegin % Alignment;
       if (TgtPadding) {
@@ -638,23 +665,23 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     void *PointerHstPtrBegin, *PointerTgtPtrBegin;
     TargetPointerResultTy PointerTpr;
     bool IsHostPtr = false;
-    bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
+    bool IsImplicit = ArgType & OMP_TGT_MAPTYPE_IMPLICIT;
     // Force the creation of a device side copy of the data when:
     // a close map modifier was associated with a map that contained a to.
-    bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
-    bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
-    bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;
+    bool HasCloseModifier = ArgType & OMP_TGT_MAPTYPE_CLOSE;
+    bool HasPresentModifier = ArgType & OMP_TGT_MAPTYPE_PRESENT;
+    bool HasHoldModifier = ArgType & OMP_TGT_MAPTYPE_OMPX_HOLD;
     // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
     // have reached this point via __tgt_target_data_begin and not __tgt_target
     // then no argument is marked as TARGET_PARAM ("omp target data map" is not
     // associated with a target region, so there are no target parameters). This
     // may be considered a hack, we could revise the scheme in the future.
     bool UpdateRef =
-        !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0);
+        !(ArgType & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0);
 
     DeviceTy::HDTTMapAccessorTy HDTTMap =
         Device.HostDataToTargetMap.getExclusiveAccessor();
-    if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
+    if (ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       DP("Has a pointer entry: \n");
       // Base is address of pointer.
       //
@@ -696,8 +723,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
           (!FromMapper || I != 0); // subsequently update ref count of pointee
     }
 
-    const bool HasFlagTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;
-    const bool HasFlagAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
+    const bool HasFlagTo = ArgType & OMP_TGT_MAPTYPE_TO;
+    const bool HasFlagAlways = ArgType & OMP_TGT_MAPTYPE_ALWAYS;
     // Note that HDTTMap will be released in getTargetPointer.
     auto TPR = Device.getTargetPointer(
         HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName,
@@ -717,14 +744,14 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
        " - is%s new\n",
        DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
 
-    if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
+    if (ArgType & OMP_TGT_MAPTYPE_RETURN_PARAM) {
       uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
       void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
       DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
       ArgsBase[I] = TgtPtrBase;
     }
 
-    if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
+    if (ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
 
       uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
       void *ExpectedTgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
@@ -752,8 +779,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     }
 
     // Check if variable can be used on the device:
-    bool IsStructMember = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
-    if (getInfoLevel() & OMP_INFOTYPE_EMPTY_MAPPING && ArgTypes[I] != 0 &&
+    bool IsStructMember = ArgType & OMP_TGT_MAPTYPE_MEMBER_OF;
+    if (getInfoLevel() & OMP_INFOTYPE_EMPTY_MAPPING && ArgType != 0 &&
         !IsStructMember && !IsImplicit && !TPR.isPresent() &&
         !TPR.isContained() && !TPR.isHostPointer())
       INFO(OMP_INFOTYPE_EMPTY_MAPPING, Device.DeviceID,
diff --git a/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp b/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp
new file mode 100644
index 000000000000000..4c9d5c7a23e8162
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp
@@ -0,0 +1,114 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+#include <stdio.h>
+#include <stdlib.h>
+
+struct Descriptor {
+  int *datum;
+  long int x;
+  int *more_datum;
+  int xi;
+  int val_datum, val_more_datum;
+  long int arr[1][30];
+  int val_arr;
+};
+
+int main() {
+  Descriptor dat = Descriptor();
+  dat.datum = (int *)malloc(sizeof(int) * 10);
+  dat.more_datum = (int *)malloc(sizeof(int) * 20);
+  dat.xi = 3;
+  dat.arr[0][0] = 1;
+
+  dat.datum[7] = 7;
+  dat.more_datum[17] = 17;
+
+  /// The struct is mapped with type 0x0 when the pointer fields are mapped.
+  /// The struct is also map explicitely by the user. The second mapping by
+  /// the user must not overwrite the mapping set up for the pointer fields
+  /// when mapping the struct happens after the mapping of the pointers.
+
+  // clang-format off
+  // CHECK: Libomptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown
+  // CHECK: Libomptarget --> Entry  1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown
+  // CHECK: Libomptarget --> Entry  2: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown
+  // CHECK: Libomptarget --> Entry  3: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown
+  // clang-format on
+
+  /// The struct will be mapped in the same order as the above entries.
+
+  /// First argument is the struct itself and it will be mapped once.
+
+  // clang-format off
+  // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[DAT_HST_PTR_BASE]], Size=288)...
+  // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 288 with host pointer [[DAT_HST_PTR_BASE]].
+  // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[DAT_HST_PTR_BASE]], HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtAllocBegin=[[DAT_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1, HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]])
+  // clang-format on
+
+  /// Second argument is dat.datum:
+  // clang-format off
+  // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[DATUM_HST_PTR_BASE]], Size=40)...
+  // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 40 with host pointer [[DATUM_HST_PTR_BASE]].
+  // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[DATUM_HST_PTR_BASE]], HstPtrBegin=[[DATUM_HST_PTR_BASE]], TgtAllocBegin=[[DATUM_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DATUM_DEVICE_PTR_BASE]], Size=40, DynRefCount=1, HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]])
+  // clang-format on
+
+  /// Third argument is dat.more_datum:
+  // clang-format off
+  // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], Size=80)...
+  // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 80 with host pointer [[MORE_DATUM_HST_PTR_BEGIN]].
+  // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[MORE_DATUM_HST_PTR_BEGIN]], HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], TgtAllocBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN:0x.*]], TgtPtrBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN]], Size=80, DynRefCount=1, HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]])
+  // clang-format on
+
+#pragma omp target enter data map(to : dat.datum[ : 10])                       \
+    map(to : dat.more_datum[ : 20]) map(to : dat)
+
+  /// Checks induced by having a target region:
+  // clang-format off
+  // CHECK: Libomptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown
+  // CHECK: Libomptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown
+  // CHECK: Libomptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]]
+  // clang-format on
+
+#pragma omp target
+  {
+    dat.xi = 4;
+    dat.datum[7]++;
+    dat.more_datum[17]++;
+    dat.val_datum = dat.datum[7];
+    dat.val_more_datum = dat.more_datum[17];
+    dat.datum[dat.arr[0][0]] = dat.xi;
+    dat.val_arr = dat.datum[dat.arr[0][0]];
+  }
+
+  /// Post-target region checks:
+  // clang-format off
+  // CHECK: Libomptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=0
+  // clang-format on
+
+#pragma omp target exit data map(from : dat)
+
+  /// Target data end checks:
+  // clang-format off
+  // CHECK: Libomptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+  // CHECK: Libomptarget --> Moving 288 bytes (tgt:[[DAT_DEVICE_PTR_BASE]]) -> (hst:[[DAT_HST_PTR_BASE]])
+  // clang-format on
+
+  // CHECK: dat.xi = 4
+  // CHECK: dat.val_datum = 8
+  // CHECK: dat.val_more_datum = 18
+  // CHECK: dat.datum[dat.arr[0][0]] = 0
+  // CHECK: dat.val_arr = 4
+
+  printf("dat.xi = %d\n", dat.xi);
+  printf("dat.val_datum = %d\n", dat.val_datum);
+  printf("dat.val_more_datum = %d\n", dat.val_more_datum);
+  printf("dat.datum[dat.arr[0][0]] = %d\n", dat.datum[dat.arr[0][0]]);
+  printf("dat.val_arr = %d\n", dat.val_arr);
+
+  return 0;
+}



More information about the Openmp-commits mailing list