[Offload]: Skip copying of unused kernel-mapped data

Author: pradt2

offload/include/OpenMP/Mapping.h

@@ -386,6 +386,13 @@ struct LookupResult {
   LookupResult() : Flags({0, 0, 0}), TPR() {}
 
   TargetPointerResultTy TPR;
+
+  bool isEmpty() const {
+      bool IsEmpty = Flags.IsContained == 0
+              & Flags.ExtendsBefore == 0
+              & Flags.ExtendsAfter == 0;
+      return IsEmpty;
+  }
 };
 
 // This structure stores information of a mapped memory region.

offload/include/Shared/Debug.h

@@ -58,6 +58,8 @@ enum OpenMPInfoType : uint32_t {
   OMP_INFOTYPE_DATA_TRANSFER = 0x0020,
   // Print whenever data does not have a viable device counterpart.
   OMP_INFOTYPE_EMPTY_MAPPING = 0x0040,
+  // Print whenever data does not need to be transferred
+  OMP_INFOTYPE_REDUNDANT_TRANSFER = 0x0080,
   // Enable every flag.
   OMP_INFOTYPE_ALL = 0xffffffff,
 };

offload/libomptarget/omptarget.cpp

@@ -1197,30 +1197,66 @@ class PrivateArgumentManagerTy {
   }
 };
 
-static std::unique_ptr<int64_t[]> maskIgnorableMappings(int64_t DeviceId, int32_t ArgNum, int64_t *ArgTypes,
-                                                        int64_t *ArgSizes, map_var_info_t *ArgNames) {
+/// Try to find redundant mappings associated with a kernel launch,
+/// and provide a masked version of the kernel argument types that
+/// avoid redundant to data transfers between the host and device.
+static std::unique_ptr<int64_t[]> maskRedundantTransfers(DeviceTy &Device, int32_t ArgNum,
+                                                         int64_t *ArgTypes, int64_t *ArgSizes,
+                                                         map_var_info_t *ArgNames, void **ArgPtrs,
+                                                         void **ArgMappers) {
     std::unique_ptr<int64_t[]> ArgTypesOverride = std::make_unique<int64_t[]>(ArgNum);
 
+    MappingInfoTy &MappingInfo = Device.getMappingInfo();
+    MappingInfoTy::HDTTMapAccessorTy HDTTMap = MappingInfo
+            .HostDataToTargetMap.getExclusiveAccessor();
+
+    int64_t UnusedArgs = 0;
+
     for (int32_t I = 0; I < ArgNum; ++I) {
-        bool IsTargetParam = ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM;
+        tgt_map_type ArgType = (tgt_map_type) ArgTypes[I];
+
+        // Check for unused implicit mappings
+        bool IsArgUnused = ArgType == OMP_TGT_MAPTYPE_NONE;
+
+        // Check for unused `map(buf[0:size])` mappings
+        IsArgUnused |= ArgType == OMP_TGT_MAPTYPE_FROM
+                || ArgType == OMP_TGT_MAPTYPE_TO
+                || ArgType == (OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO);
+
+        // Check for unused `map(wrapper.buf[0:size])` mappings
+        IsArgUnused |= UnusedArgs == ArgNum - 1 && ArgType & OMP_TGT_MAPTYPE_MEMBER_OF
+                && ((ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == OMP_TGT_MAPTYPE_PTR_AND_OBJ
+                || (ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == (OMP_TGT_MAPTYPE_PTR_AND_OBJ | OMP_TGT_MAPTYPE_TO)
+                || (ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == (OMP_TGT_MAPTYPE_PTR_AND_OBJ | OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO));
 
-        bool IsMapTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;
-        if (IsTargetParam || !IsMapTo) {
+        bool IsExistingMapping = !MappingInfo.lookupMapping(HDTTMap, ArgPtrs[I], ArgSizes[I]).isEmpty();
+
+        bool IsCustomMapped = ArgMappers && ArgMappers[I];
+
+        if (IsExistingMapping | IsCustomMapped | !IsArgUnused) {
             ArgTypesOverride[I] = ArgTypes[I];
             continue;
         }
 
-        bool IsMapFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
-        const char *Type = IsMapFrom ? "tofrom" : "to";
+        const std::string Name = ArgNames && ArgNames[I] ?
+                                 getNameFromMapping(ArgNames[I]) : std::string("unknown");
+
+        bool IsArgFrom = ArgType & OMP_TGT_MAPTYPE_FROM;
+        bool IsArgTo = ArgType & OMP_TGT_MAPTYPE_TO;
+
+        const char *Type = IsArgFrom && IsArgTo ? "tofrom"
+                               : IsArgFrom ? "from"
+                               : IsArgTo ? "to"
+                               : "unknown";
 
-        // Optimisation: A 'to' or 'tofrom' mapping is not
-        // used by the kernel. Change its type such that
-        // no new mapping is created, but any existing
-        // mapping has its counter decremented.
-        INFO(OMP_INFOTYPE_ALL, DeviceId, "%s(%s)[%" PRId64 "] %s\n", Type,
-             getNameFromMapping(ArgNames[I]).c_str(), ArgSizes[I], "is not used and will not be copied");
+        // Optimisation:
+        // A new mapping is not used by the kernel.
+        // Change the type such that no data is transferred to and/or from the device.
+        INFO(OMP_INFOTYPE_REDUNDANT_TRANSFER, Device.DeviceID, "%s(%s)[%" PRId64 "] %s\n", Type,
+             Name.c_str(), ArgSizes[I], "is not used and will not be copied");
 
         ArgTypesOverride[I] = ArgTypes[I] & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM);
+        UnusedArgs++;
     }
 
     return ArgTypesOverride;
@@ -1448,8 +1484,8 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
   int Ret = OFFLOAD_SUCCESS;
 
     std::unique_ptr<int64_t[]> ArgTypesOverride =
-            maskIgnorableMappings(DeviceId, NumClangLaunchArgs, KernelArgs.ArgTypes,
-                                  KernelArgs.ArgSizes, KernelArgs.ArgNames);
+            maskRedundantTransfers(Device, NumClangLaunchArgs, KernelArgs.ArgTypes,
+                                  KernelArgs.ArgSizes, KernelArgs.ArgNames, KernelArgs.ArgPtrs, KernelArgs.ArgMappers);
 
   if (NumClangLaunchArgs) {
     // Process data, such as data mapping, before launching the kernel

offload/test/mapping/skip_transfers.cpp

@@ -0,0 +1,72 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic
+// RUN: env LIBOMPTARGET_INFO=128 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// REQUIRES: gpu
+// clang-format on
+
+int main() {
+    float DataStack = 0;
+
+    // CHECK-NOT: omptarget device 0 info: from(unknown)[4] is not used and will not be copied
+    #pragma omp target map(from: DataStack)
+    {
+        DataStack = 1;
+    }
+
+    // CHECK-NOT: omptarget device 0 info: to(unknown)[4] is not used and will not be copied
+    #pragma omp target map(always to: DataStack)
+    ;
+
+    // CHECK: omptarget device 0 info: tofrom(unknown)[4] is not used and will not be copied
+    #pragma omp target map(tofrom: DataStack)
+    ;
+
+    int Size = 16;
+    double *Data = new double[Size];
+
+    // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[8] is not used and will not be copied
+    #pragma omp target map(tofrom: Data[0:1])
+    {
+        Data[0] = 1;
+    }
+
+    // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[16] is not used and will not be copied
+    #pragma omp target map(always tofrom: Data[0:2])
+    ;
+
+    // CHECK: omptarget device 0 info: from(unknown)[24] is not used and will not be copied
+    #pragma omp target map(from: Data[0:3])
+    ;
+
+    // CHECK: omptarget device 0 info: to(unknown)[24] is not used and will not be copied
+    #pragma omp target map(to: Data[0:3])
+    ;
+
+    // CHECK: omptarget device 0 info: tofrom(unknown)[32] is not used and will not be copied
+    #pragma omp target map(tofrom: Data[0:4])
+    ;
+
+    // CHECK-NOT: omptarget device 0 info: to(unknown)[40] is not used and will not be copied
+    #pragma omp target map(to: Data[0:5])
+    {
+        #pragma omp teams
+        Data[0] = 1;
+    }
+
+    struct {
+      double *Data;
+    } Wrapper { .Data = Data };
+
+    // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[48] is not used and will not be copied
+    #pragma omp target map(tofrom: Wrapper.Data[0:6])
+    {
+        Wrapper.Data[0] = 1;
+    }
+
+    // CHECK: omptarget device 0 info: unknown(unknown)[8] is not used and will not be copied
+    // CHECK: omptarget device 0 info: tofrom(unknown)[56] is not used and will not be copied
+    #pragma omp target map(tofrom: Wrapper.Data[0:7])
+    ;
+}