[Openmp-commits] [openmp] bb0166d - [libomptarget] Update device pointer only if needed

[George Rokos via Openmp-commits]

Author: George Rokos
Date: 2021-07-13T04:18:55-07:00
New Revision: bb0166dc72791e2cefdb0c8dc9e495ea0555357b

URL: https://github.com/llvm/llvm-project/commit/bb0166dc72791e2cefdb0c8dc9e495ea0555357b
DIFF: https://github.com/llvm/llvm-project/commit/bb0166dc72791e2cefdb0c8dc9e495ea0555357b.diff

LOG: [libomptarget] Update device pointer only if needed

Currently, libomptarget will always perform a host-to-device memory transfer in
order to update the device pointer of a PTR_AND_OBJ entry. This is not always
necessary because the device pointer may have been set to the correct pointee
address already, so we can eliminate the redundant memory transfer.

Added: 
    openmp/libomptarget/test/mapping/device_ptr_update.c

Modified: 
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 80cb9381ffbb5..0255f94599ee1 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -556,22 +556,35 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
-      DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n",
-         DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
+      // Check whether we need to update the pointer on the device
+      bool UpdateDevPtr = false;
+
       uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
-      void *&TgtPtrBase = AsyncInfo.getVoidPtrLocation();
-      TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
-      int rt = Device.submitData(PointerTgtPtrBegin, &TgtPtrBase,
-                                 sizeof(void *), AsyncInfo);
-      if (rt != OFFLOAD_SUCCESS) {
-        REPORT("Copying data to device failed.\n");
-        return OFFLOAD_FAIL;
-      }
-      // create shadow pointers for this entry
+      void *ExpectedTgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
+
       Device.ShadowMtx.lock();
-      Device.ShadowPtrMap[Pointer_HstPtrBegin] = {
-          HstPtrBase, PointerTgtPtrBegin, TgtPtrBase};
+      auto Entry = Device.ShadowPtrMap.find(Pointer_HstPtrBegin);
+      // If this pointer is not in the map we need to insert it.
+      if (Entry == Device.ShadowPtrMap.end()) {
+        // create shadow pointers for this entry
+        Device.ShadowPtrMap[Pointer_HstPtrBegin] = {
+            HstPtrBase, PointerTgtPtrBegin, ExpectedTgtPtrBase};
+        UpdateDevPtr = true;
+      }
       Device.ShadowMtx.unlock();
+
+      if (UpdateDevPtr) {
+        DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n",
+           DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
+        void *&TgtPtrBase = AsyncInfo.getVoidPtrLocation();
+        TgtPtrBase = ExpectedTgtPtrBase;
+        int rt = Device.submitData(PointerTgtPtrBegin, &TgtPtrBase,
+                                   sizeof(void *), AsyncInfo);
+        if (rt != OFFLOAD_SUCCESS) {
+          REPORT("Copying data to device failed.\n");
+          return OFFLOAD_FAIL;
+        }
+      }
     }
   }
 

diff  --git a/openmp/libomptarget/test/mapping/device_ptr_update.c b/openmp/libomptarget/test/mapping/device_ptr_update.c
new file mode 100644
index 0000000000000..e3dd12f3d4340
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/device_ptr_update.c
@@ -0,0 +1,44 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// REQUIRES: libomptarget-debug
+
+#include <stdio.h>
+
+struct S {
+  int *p;
+};
+
+int main(void) {
+  int A[10];
+  struct S s1;
+
+  s1.p = A;
+
+  // DEBUG: Update pointer ([[DEV_PTR:0x[^ ]+]]) -> {{\[}}[[DEV_OBJ_A:0x[^ ]+]]{{\]}}
+  #pragma omp target enter data map(alloc : s1.p [0:10])
+
+  // DEBUG-NOT: Update pointer ([[DEV_PTR]]) -> {{\[}}[[DEV_OBJ_A]]{{\]}}
+  #pragma omp target map(alloc : s1.p [0:10])
+  {
+    for (int i = 0; i < 10; ++i)
+      s1.p[i] = i;
+  }
+
+  #pragma omp target exit data map(from : s1.p [0:10])
+
+  int fail_A = 0;
+  for (int i = 0; i < 10; ++i) {
+    if (A[i] != i) {
+      fail_A = 1;
+      break;
+    }
+  }
+
+  // CHECK-NOT: Test A failed
+  if (fail_A) {
+    printf("Test A failed\n");
+  }
+
+  return fail_A;
+}