[OpenMP] Fix `omp target update` for array extension

[jdenny]

Summary

OpenMP TR8 sec. 2.15.6 "target update Construct", p. 183, L3-4 states:

If the corresponding list item is not present in the device data
environment and there is no present modifier in the clause, then no
assignment occurs to or from the original list item.

L10-11 states:

If a present modifier appears in the clause and the corresponding
list item is not present in the device data environment then an
error occurs and the program termintates.

(OpenMP 5.0 also has the first passage but without mention of the
present modifier of course.)

In both passages, I assume "is not present" includes the case of
partially but not entirely present. However, without this patch, the
target update directive misbehaves in this case both with and without
the present modifier. For example:

#pragma omp target enter data map(to:arr[0:3])
#pragma omp target update to(arr[0:5]) // might fail on data transfer
#pragma omp target update to(present:arr[0:5]) // might fail on data transfer

The problem is that DeviceTy::getTgtPtrBegin does not return a null
pointer in that case, so target_data_update sees the data as fully
present, and the data transfer then might fail depending on the target
device. However, without the present modifier, there should never be
a failure. Moreover, with the present modifier, there should always
be a failure, and the diagnostic should mention the present modifier.

This patch fixes DeviceTy::getTgtPtrBegin to return null when
target_data_update is the caller. I'm wondering if it should do the
same for more callers.

[grokos]

I suppose the same must be applied to all cases where present is used. Without a modification like this, present just confirms that the starting address of the object is already mapped, although a size may be specified in the map clause.

#pragma omp target map(present: p[0:10]) // present will succeed even if only p[0] is mapped and p[1]-p[9] are not

So the question here is: should present apply to the size as well as the begin address? I would say yes.

[jdenny]

Thanks for the quick review.

In D85246#2194712, @grokos wrote:

I suppose the same must be applied to all cases where present is used.

Agreed. Fortunately, when present is instead a map type modifier, DeviceTy::getOrAllocTgtPtr is called instead of DeviceTy::getTgtPtrBegin and, as far as I know, it already returns a null pointer for this case. However, I see now that we get the general failure that happens regardless of the present modifier, and we don't see the present-specific diagnostic. I suppose that ought to change.

Without a modification like this, present just confirms that the starting address of the object is already mapped, although a size may be specified in the map clause.

It doesn't even confirm the start address exactly. That is, in the extends-before case, the starting address might be before the array section that was already mapped. Without this patch, the present motion modifier permits that case too.

#pragma omp target map(present: p[0:10]) // present will succeed even if only p[0] is mapped and p[1]-p[9] are not

So the question here is: should present apply to the size as well as the begin address? I would say yes.

I agree, and it seems to already except for the diagnostic.

[jdoerfert]

Maybe change the commit title to OpenMP to confuse less people.

[jdenny]

In D85246#2194893, @jdoerfert wrote:

Maybe change the commit title to OpenMP to confuse less people.

Thanks for catching that. Done.

[jdenny]

This patch fixes DeviceTy::getTgtPtrBegin to return null when
target_data_update is the caller. I'm wondering if it should do the
same for more callers.

Here's an example of what I'm thinking here:

#pragma omp target data map(tofrom:arr[0:20])
{
  #pragma omp target exit data map(delete:arr[0:20])
  #pragma omp target enter data map(alloc:arr[10:2])
}

The transfer at the end of the target data region currently fails. According to my read of the spec, the transfer shouldn't even be attempted because arr[0:20] isn't (fully) present. If DeviceTy::getTgtPtrBegin were to return null for that case too, then this problem would go away.

openmp/libomptarget/src/device.h

struct DeviceTy {
   uint64_t getMapEntryRefCnt(void *HstPtrBegin);
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
                          bool &IsNew, bool &IsHostPtr, bool IsImplicit,
                          bool UpdateRefCount, bool HasCloseModifier,
                          bool HasPresentModifier);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-      bool UpdateRefCount, bool &IsHostPtr);
+                       bool UpdateRefCount, bool &IsHostPtr,
+                       bool MustContain = false);
   int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
                     bool HasCloseModifier = false);
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);
 
   // calls to RTL
   int32_t initOnce();
   __tgt_target_table *load_binary(void *Img);

openmp/libomptarget/src/device.cpp

void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
   DataMapMtx.unlock();
   return rc;
 }
 
 // Used by targetDataBegin, targetDataEnd, target_data_update and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from targetDataEnd.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-    bool UpdateRefCount, bool &IsHostPtr) {
+                               bool UpdateRefCount, bool &IsHostPtr,
+                               bool MustContain) {
   void *rc = NULL;
   IsHostPtr = false;
   IsLast = false;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
 
-  if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
+  if (lr.Flags.IsContained ||
+      (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
     auto &HT = *lr.Entry;
     IsLast = HT.getRefCount() == 1;
 
     if (!IsLast && UpdateRefCount)
       HT.decRefCount();
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "

openmp/libomptarget/src/omptarget.cpp

if (arg_mappers && arg_mappers[i]) {
 
       // Skip the rest of this function, continue to the next argument.
       continue;
     }
 
     void *HstPtrBegin = args[i];
     int64_t MapSize = arg_sizes[i];
     bool IsLast, IsHostPtr;
-    void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast,
-        false, IsHostPtr);
+    void *TgtPtrBegin = Device.getTgtPtrBegin(
+        HstPtrBegin, MapSize, IsLast, false, IsHostPtr, /*MustContain=*/true);
     if (!TgtPtrBegin) {
       DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
       if (arg_types[i] & OMP_TGT_MAPTYPE_PRESENT) {
         MESSAGE("device mapping required by 'present' motion modifier does not "
                 "exist for host address " DPxMOD " (%ld bytes)",
                 DPxPTR(HstPtrBegin), MapSize);
         return OFFLOAD_FAIL;
       }

openmp/libomptarget/test/mapping/present/target_update_array_extension.c

+// --------------------------------------------------
+// Check 'to' and extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'from' and extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'to' and extends after
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'from' and extends after
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// END.
+
+#include <stdio.h>
+
+#define BEFORE 0
+#define AFTER  1
+
+#if EXTENDS == BEFORE
+# define SMALL 2:3
+# define LARGE 0:5
+#elif EXTENDS == AFTER
+# define SMALL 0:3
+# define LARGE 0:5
+#else
+# error EXTENDS undefined
+#endif
+
+int main() {
+  int arr[5];
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+  fprintf(stderr, "addr=%p, size=%ld\n", arr, sizeof arr);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: arr[LARGE])
+  {
+#pragma omp target update CLAUSE(present: arr[SMALL])
+  }
+
+  // CHECK: arr is present
+  fprintf(stderr, "arr is present\n");
+
+  // CHECK: Libomptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+  // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target data map(alloc: arr[SMALL])
+  {
+#pragma omp target update CLAUSE(present: arr[LARGE])
+  }
+
+  // CHECK-NOT: arr is present
+  fprintf(stderr, "arr is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/target_update_array_extension.c

+// --------------------------------------------------
+// Check 'to' and extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'from' and extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'to' and extends after
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -DCLAUSE=to -DEXTENDS=AFTER
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'from' and extends after
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
+// RUN:   -DCLAUSE=from -DEXTENDS=AFTER
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// END.
+
+#include <stdio.h>
+
+#define BEFORE 0
+#define AFTER  1
+
+#if EXTENDS == BEFORE
+# define SMALL 2:3
+# define LARGE 0:5
+#elif EXTENDS == AFTER
+# define SMALL 0:3
+# define LARGE 0:5
+#else
+# error EXTENDS undefined
+#endif
+
+int main() {
+  int arr[5];
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: arr[LARGE])
+  {
+#pragma omp target update CLAUSE(arr[SMALL])
+  }
+
+  // CHECK: success
+  fprintf(stderr, "success\n");
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: arr[SMALL])
+  {
+#pragma omp target update CLAUSE(arr[LARGE])
+  }
+
+  // CHECK: success
+  fprintf(stderr, "success\n");
+
+  return 0;
+}