[libomptarget] Update device pointer only if needed

[grokos]

Summary

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.

The attached test is C code emulating what happens when mapping Fortran dope vectors.

[tianshilei1992]

LGTM. Thanks. Please fix clang-format issues before commit.

[jdenny]

openmp/libomptarget/test/mapping/device_ptr_update.c

20

Assume we insert s1.p = B here. Without this patch, the target directive below then changes s1.p on the device, but it doesn't with this patch, right?

However, I don't think OpenMP permits that line to be inserted in the case of C/C++:

A pointer that has a corresponding attached pointer must not be modified for the duration of the lifetime of the list item to which the corresponding pointer is attached in the device data environment.

This patch can then be considered as purely an optimization, right?

[grokos]

openmp/libomptarget/test/mapping/device_ptr_update.c

20

Right. Since the pointer on the host is not allowed to be reassigned to another object while the PTR_AND_OBJ pair is mapped on the device, we can take advantage of this restriction and optimize away all redundant device pointer updates.

[grokos]

Closed by commit rGbb0166dc7279: [libomptarget] Update device pointer only if needed (authored by grokos).

[jdenny]

openmp/libomptarget/test/mapping/device_ptr_update.c

20

Thanks.

openmp/libomptarget/src/omptarget.cpp

if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
         if (rt != OFFLOAD_SUCCESS) {
           REPORT("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
+      // Check whether we need to update the pointer on the device
+      bool UpdateDevPtr = false;
+
+      uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
+      void *ExpectedTgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
+
+      Device.ShadowMtx.lock();
+      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));
-      uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
         void *&TgtPtrBase = AsyncInfo.getVoidPtrLocation();
-      TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
+        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;
         }
-      // create shadow pointers for this entry
-      Device.ShadowMtx.lock();
-      Device.ShadowPtrMap[Pointer_HstPtrBegin] = {
-          HstPtrBase, PointerTgtPtrBegin, TgtPtrBase};
-      Device.ShadowMtx.unlock();
+      }
     }
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 namespace {
 /// This structure contains information to deallocate a target pointer, aka.

openmp/libomptarget/test/mapping/device_ptr_update.c

+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG
+// 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;
+}