[libomptarget] Implement pointer lookup as 5.1 spec.

[ye-luo]

Summary

As described in 5.1 spec
2.21.7.2 Pointer Initialization for Device Data Environments

[ye-luo]

Two tests array_section_use_device_ptr.c and array_section_implicit_capture.c are ported from https://reviews.llvm.org/D117997

[RaviNarayanaswamy]

LGTM

openmp/libomptarget/src/device.cpp

LookupResult DeviceTy::lookupMapping(HDTTMapAccessorTy &HDTTMap,
 
   DP("Looking up mapping(HstPtrBegin=" DPxMOD ", Size=%" PRId64 ")...\n",
      DPxPTR(hp), Size);
 
   if (HDTTMap->empty())
     return lr;
 
   auto upper = HDTTMap->upper_bound(hp);
+
+  if (Size == 0) {
+    // specification v5.1 Pointer Initialization for Device Data Environments
+    // upper_bound satisfies
+    //   std::prev(upper)->HDTT.HstPtrBegin <= hp < upper->HDTT.HstPtrBegin
+    if (upper != HDTTMap->begin()) {
+      lr.Entry = std::prev(upper)->HDTT;
+      auto &HT = *lr.Entry;
+      // the left side of extended address range is satisified.
+      // hp >= HT.HstPtrBegin || hp >= HT.HstPtrBase
+      lr.Flags.IsContained = hp < HT.HstPtrEnd || hp < HT.HstPtrBase;
+    }
+
+    if (!lr.Flags.IsContained && upper != HDTTMap->end()) {
+      lr.Entry = upper->HDTT;
+      auto &HT = *lr.Entry;
+      // the right side of extended address range is satisified.
+      // hp < HT.HstPtrEnd || hp < HT.HstPtrBase
+      lr.Flags.IsContained = hp >= HT.HstPtrBase;
+    }
+  } else {
     // check the left bin
     if (upper != HDTTMap->begin()) {
       lr.Entry = std::prev(upper)->HDTT;
       auto &HT = *lr.Entry;
       // Is it contained?
       lr.Flags.IsContained = hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd &&
                              (hp + Size) <= HT.HstPtrEnd;
       // Does it extend beyond the mapped region?
       lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd;
     }
 
     // check the right bin
     if (!(lr.Flags.IsContained || lr.Flags.ExtendsAfter) &&
         upper != HDTTMap->end()) {
       lr.Entry = upper->HDTT;
       auto &HT = *lr.Entry;
       // Does it extend into an already mapped region?
       lr.Flags.ExtendsBefore =
           hp < HT.HstPtrBegin && (hp + Size) > HT.HstPtrBegin;
       // Does it extend beyond the mapped region?
       lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd;
     }
 
     if (lr.Flags.ExtendsBefore) {
       DP("WARNING: Pointer is not mapped but section extends into already "
          "mapped data\n");
     }
     if (lr.Flags.ExtendsAfter) {
       DP("WARNING: Pointer is already mapped but section extends beyond mapped "
          "region\n");
     }
+  }
 
   return lr;
 }
 
 TargetPointerResultTy DeviceTy::getTargetPointer(
     void *HstPtrBegin, void *HstPtrBase, int64_t Size,
     map_var_info_t HstPtrName, bool HasFlagTo, bool HasFlagAlways,
     bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
if (LR.Flags.IsContained ||
     uintptr_t Ptr = (uintptr_t)allocData(Size, HstPtrBegin);
     Entry = HDTTMap
                 ->emplace(new HostDataToTargetTy(
                     (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
                     (uintptr_t)HstPtrBegin + Size, Ptr, HasHoldModifier,
                     HstPtrName))
                 .first->HDTT;
     INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
-         "Creating new map entry with "
-         "HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
+         "Creating new map entry with HstPtrBase= " DPxMOD
+         ", HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
          "DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
-         DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
+         DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
          Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(),
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     TargetPointer = (void *)Ptr;
   }
 
   // If the target pointer is valid, and we need to transfer data, issue the
   // data transfer.
   if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways)) {

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1024
+#define FROM 64
+#define LENGTH 128
+
+int main() {
+  float *A = (float *)malloc(N * sizeof(float));
+  float *B = (float *)malloc(N * sizeof(float));
+  float *C = (float *)malloc(N * sizeof(float));
+
+  for (int i = 0; i < N; i++) {
+    C[i] = 0.0;
+  }
+
+  for (int i = 0; i < N; i++) {
+    A[i] = i;
+    B[i] = 2 * i;
+  }
+
+#pragma omp target enter data map(to : A [FROM:LENGTH], B [FROM:LENGTH])
+#pragma omp target enter data map(alloc : C [FROM:LENGTH])
+
+// A, B and C have been mapped starting at index FROM, but inside the kernel
+// they are captured implicitly so the library must look them up using their
+// base address.
+#pragma omp target
+  {
+    for (int i = FROM; i < FROM + LENGTH; i++) {
+      C[i] = A[i] + B[i];
+    }
+  }
+
+#pragma omp target exit data map(from : C [FROM:LENGTH])
+#pragma omp target exit data map(delete : A [FROM:LENGTH], B [FROM:LENGTH])
+
+  int errors = 0;
+  for (int i = FROM; i < FROM + LENGTH; i++)
+    if (C[i] != A[i] + B[i])
+      ++errors;
+
+  // CHECK: Success
+  if (errors)
+    fprintf(stderr, "Failure\n");
+  else
+    fprintf(stderr, "Success\n");
+
+  free(A);
+  free(B);
+  free(C);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/array_section_use_device_ptr.c

+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1024
+#define FROM 64
+#define LENGTH 128
+
+int main() {
+  float *A = (float *)malloc(N * sizeof(float));
+
+#pragma omp target enter data map(to : A [FROM:LENGTH])
+
+  // A, has been mapped starting at index FROM, but inside the use_device_ptr
+  // clause it is captured by base so the library must look it up using the
+  // base address.
+
+  float *A_dev = NULL;
+#pragma omp target data use_device_ptr(A)
+  { A_dev = A; }
+#pragma omp target exit data map(delete : A [FROM:LENGTH])
+
+  // CHECK: Success
+  if (A_dev == NULL || A_dev == A)
+    fprintf(stderr, "Failure\n");
+  else
+    fprintf(stderr, "Success\n");
+
+  free(A);
+
+  return 0;
+}