[Openmp-commits] [PATCH] D123892: [libomptarget] add a test mapping three chunks of the same array

[Ye Luo via Phabricator via Openmp-commits]

ye-luo created this revision.
Herald added a project: All.
ye-luo requested review of this revision.
Herald added a project: OpenMP.
Herald added a subscriber: openmp-commits.

Add an interesting test. As long as only one chunk is used in a target region and explicit map calls out the desired chunk, mapping non-overlapping chunks of an array should work.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D123892

Files:
  openmp/libomptarget/src/device.cpp
  openmp/libomptarget/test/mapping/three_chunks_of_one_array.c


Index: openmp/libomptarget/test/mapping/three_chunks_of_one_array.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/three_chunks_of_one_array.c
@@ -0,0 +1,32 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+  int *mem = (int *)malloc(32 * sizeof(int));
+  mem[0] = 1;
+  mem[5] = 11;
+  mem[20] = 21;
+#pragma omp target data map(to : mem [20:10], mem [5:5]) map(to : mem[0])
+  {
+#pragma omp parallel
+#pragma omp single
+    {
+#pragma omp target map(always, from : mem[0]) nowait
+      { mem[0] = 2; }
+#pragma omp target map(always, from : mem [5:5]) nowait
+      { mem[5] = 12; }
+#pragma omp target map(always, from : mem [20:10]) nowait
+      { mem[20] = 22; }
+    }
+    // CHECK: mem[0] 2
+    printf("mem[0] %d\n", mem[0]);
+    // CHECK: mem[5] 12
+    printf("mem[5] %d\n", mem[5]);
+    // CHECK: mem[20] 22
+    printf("mem[20] %d\n", mem[20]);
+  }
+
+  return 0;
+}
Index: openmp/libomptarget/src/device.cpp
===================================================================
--- openmp/libomptarget/src/device.cpp
+++ openmp/libomptarget/src/device.cpp
@@ -245,11 +245,12 @@
     const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
     uintptr_t Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
-         "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
-         ", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
-         (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr),
-         Size, HT.dynRefCountToStr().c_str(), DynRefCountAction,
-         HT.holdRefCountToStr().c_str(), HoldRefCountAction,
+         "Mapping exists%s with HstPtrBase=" DPxMOD ", HstPtrBegin=" DPxMOD
+         ", TgtPtrBegin=" DPxMOD ", Size=%" PRId64
+         ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
+         (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBase),
+         DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size, HT.dynRefCountToStr().c_str(),
+         DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction,
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     TargetPointer = (void *)Ptr;
   } else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D123892.423217.patch
Type: text/x-patch
Size: 2421 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20220416/d9d1bad5/attachment-0001.bin>