[llvm] [OpenMP] Fix crash with duplicate mapping on target directive (PR #146136)

[Julian Brown via llvm-commits]

https://github.com/jtb20 updated https://github.com/llvm/llvm-project/pull/146136

>From f7c5d894d32929f7d4eefc9314123f2f80dd6335 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian.brown at amd.com>
Date: Fri, 27 Jun 2025 07:08:02 -0500
Subject: [PATCH] [OpenMP] Fix crash with duplicate mapping on target directive

OpenMP allows duplicate mappings, i.e. in OpenMP 6.0, 7.9.6 "map Clause":

  "Two list items of the map clauses on the same construct must not
   share original storage unless one of the following is true: they
   are the same list item [or other omitted reasons]"

Duplicate mappings can arise as a result of user-defined mapper processing
(which I think is a separate bug, and is not addressed here), but also
in straightforward cases such as:

  #pragma omp target map(tofrom: s.mem[0:10]) map(tofrom: s.mem[0:10])

Both these cases cause crashes at runtime at present, due to an
unfortunate interaction between reference counting behaviour and shadow
pointer handling for blocks.  This is what happens:

  1. The member "s.mem" is copied to the target
  2. A shadow pointer is created, modifying the pointer on the target
  3. The member "s.mem" is copied to the target again
  4. The previous shadow pointer metadata is still present, so the
     runtime doesn't modify the target pointer a second time.

The fix is to disable step 3 if we've already done step 2 for a given
block that has the "is new" flag set.
---
 offload/libomptarget/OpenMP/Mapping.cpp       | 22 ++++++++++++++
 offload/test/mapping/duplicate_mappings_1.cpp | 27 +++++++++++++++++
 offload/test/mapping/duplicate_mappings_2.cpp | 29 +++++++++++++++++++
 3 files changed, 78 insertions(+)
 create mode 100644 offload/test/mapping/duplicate_mappings_1.cpp
 create mode 100644 offload/test/mapping/duplicate_mappings_2.cpp

diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 14f5e7dc9d19f..9b3533895f2a6 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -326,6 +326,28 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
   // data transfer.
   if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo &&
       (LR.TPR.Flags.IsNewEntry || HasFlagAlways) && Size != 0) {
+
+    // If we have something like:
+    //   #pragma omp target map(to: s.myarr[0:10]) map(to: s.myarr[0:10])
+    // then we see two "new" mappings of the struct member s.myarr here --
+    // and both have the "IsNewEntry" flag set, so trigger the copy to device
+    // below.  But, the shadow pointer is only initialised on the target for
+    // the first copy, and the second copy clobbers it.  So, this condition
+    // avoids the (second) copy here if we have already set shadow pointer info.
+    auto FailOnPtrFound = [HstPtrBegin, Size](ShadowPtrInfoTy &SP) {
+      if (SP.HstPtrAddr >= HstPtrBegin &&
+          SP.HstPtrAddr < (void *)((char *)HstPtrBegin + Size))
+        return OFFLOAD_FAIL;
+      return OFFLOAD_SUCCESS;
+    };
+    if (LR.TPR.getEntry()->foreachShadowPointerInfo(FailOnPtrFound) ==
+        OFFLOAD_FAIL) {
+      DP("Multiple new mappings of %" PRId64 " bytes detected (hst:" DPxMOD
+         ") -> (tgt:" DPxMOD ")\n",
+         Size, DPxPTR(HstPtrBegin), DPxPTR(LR.TPR.TargetPointer));
+      return std::move(LR.TPR);
+    }
+
     DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", Size,
        DPxPTR(HstPtrBegin), DPxPTR(LR.TPR.TargetPointer));
 
diff --git a/offload/test/mapping/duplicate_mappings_1.cpp b/offload/test/mapping/duplicate_mappings_1.cpp
new file mode 100644
index 0000000000000..5bed3f161b4fe
--- /dev/null
+++ b/offload/test/mapping/duplicate_mappings_1.cpp
@@ -0,0 +1,27 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic -Wno-openmp-mapping && %libomptarget-run-generic
+
+// clang-format on
+
+#include <assert.h>
+
+struct Inner {
+  int *data;
+  Inner(int size) { data = new int[size](); }
+  ~Inner() { delete[] data; }
+};
+
+struct Outer {
+  Inner i;
+  Outer() : i(10) {}
+};
+
+int main() {
+  Outer o;
+#pragma omp target map(tofrom : o.i.data[0 : 10]) map(tofrom : o.i.data[0 : 10])
+  {
+    o.i.data[0] = 42;
+  }
+  assert(o.i.data[0] == 42);
+  return 0;
+}
diff --git a/offload/test/mapping/duplicate_mappings_2.cpp b/offload/test/mapping/duplicate_mappings_2.cpp
new file mode 100644
index 0000000000000..ca8112ee72544
--- /dev/null
+++ b/offload/test/mapping/duplicate_mappings_2.cpp
@@ -0,0 +1,29 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic -Wno-openmp-mapping && %libomptarget-run-generic
+
+#include <assert.h>
+
+// clang-format on
+
+struct Inner {
+  int *data;
+  Inner(int size) { data = new int[size](); }
+  ~Inner() { delete[] data; }
+};
+#pragma omp declare mapper(Inner i) map(i, i.data[0 : 10])
+
+struct Outer {
+  Inner i;
+  Outer() : i(10) {}
+};
+#pragma omp declare mapper(Outer o) map(o, o.i)
+
+int main() {
+  Outer o;
+#pragma omp target map(tofrom : o)
+  {
+    o.i.data[0] = 42;
+  }
+  assert(o.i.data[0] == 42);
+  return 0;
+}