[libomptarget] Implement pointer lookup as 5.1 spec.

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

Reviewed By: RaviNarayanaswamy

Differential Revision: https://reviews.llvm.org/D123093

openmp/libomptarget/src/device.cpp

@@ -151,36 +151,58 @@ LookupResult DeviceTy::lookupMapping(HDTTMapAccessorTy &HDTTMap,
     return lr;
 
   auto upper = HDTTMap->upper_bound(hp);
-  // 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 (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.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");
+    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;
@@ -275,10 +297,10 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
                     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;

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

@@ -0,0 +1,58 @@
+// 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

@@ -0,0 +1,35 @@
+// 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;
+}