Skip to content

[OpenMP] Fix crash with duplicate mapping on target directive #146136

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jun 29, 2025

Conversation

jtb20
Copy link
Contributor

@jtb20 jtb20 commented Jun 27, 2025

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.

OK?

@jtb20 jtb20 requested a review from jdoerfert June 27, 2025 18:42
@llvmbot
Copy link
Member

llvmbot commented Jun 27, 2025

@llvm/pr-subscribers-openmp

@llvm/pr-subscribers-offload

Author: Julian Brown (jtb20)

Changes

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.

OK?


Full diff: https://github.com/llvm/llvm-project/pull/146136.diff

3 Files Affected:

  • (modified) offload/libomptarget/OpenMP/Mapping.cpp (+22)
  • (added) offload/test/mapping/duplicate_mappings_1.cpp (+27)
  • (added) offload/test/mapping/duplicate_mappings_2.cpp (+29)
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..ad898f80d021d
--- /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;
+}

@jtb20 jtb20 added openmp openmp:libomptarget OpenMP offload runtime labels Jun 27, 2025
@jtb20 jtb20 requested review from alexey-bataev and ro-i June 27, 2025 18:43
Copy link

github-actions bot commented Jun 27, 2025

✅ With the latest revision this PR passed the C/C++ code formatter.

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.
@jtb20 jtb20 force-pushed the duplicate-mapping-fix branch from e3005e5 to f7c5d89 Compare June 27, 2025 18:49
@jtb20
Copy link
Contributor Author

jtb20 commented Jun 27, 2025

(Generally it'd be nice if the compiler unified these kinds of mapping before the runtime sees them, but the runtime should still be able to cope with the degenerate cases rather than crashing.)

@ro-i
Copy link
Contributor

ro-i commented Jun 27, 2025

Thanks for fixing this!

I'm a little confused by what you wrote about the runtime behavior. In the code comment, you wrote:

[...] 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.

In your PR post, you mentioned that for step 4 "the previous shadow pointer metadata is still present, so the runtime doesn't modify the target pointer a second time."

So, the runtime (before the fix) didn't do the same data copy twice, but instead overwrote the previously correct shadow pointer info with the host pointer metadata leading to the target pointer point to the host memory?

@jtb20
Copy link
Contributor Author

jtb20 commented Jun 27, 2025

Thanks for fixing this!

I'm a little confused by what you wrote about the runtime behavior. In the code comment, you wrote:

[...] 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.

In your PR post, you mentioned that for step 4 "the previous shadow pointer metadata is still present, so the runtime doesn't modify the target pointer a second time."

So, the runtime (before the fix) didn't do the same data copy twice, but instead overwrote the previously correct shadow pointer info with the host pointer metadata leading to the target pointer point to the host memory?

There are three copy operations that take place, and a fourth that doesn't. So we have something like:

  • copy partial struct "s.mem" from host to target, pointing to host memory.
  • update the s.mem pointer to point to target memory instead (and create an entry in the "States->ShadowPtrInfos" set).
  • copy partial struct "s.mem" from host to target again, pointing to host memory.
  • now, addShadowPointer sees that we have an entry in States->ShadowPtrInfos, so doesn't modify the s.mem pointer again. So, we try to access the host version of the pointer on the target.

Those are just the operations corresponding to the pointer/struct member "s.mem" itself, not the separately-allocated block it points to. HTH!

Copy link
Contributor

@mjklemm mjklemm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@jtb20 jtb20 merged commit b62b58d into llvm:main Jun 29, 2025
7 checks passed
@jtb20 jtb20 deleted the duplicate-mapping-fix branch June 29, 2025 21:41
rlavaee pushed a commit to rlavaee/llvm-project that referenced this pull request Jul 1, 2025
…46136)

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.
rlavaee pushed a commit to rlavaee/llvm-project that referenced this pull request Jul 1, 2025
…46136)

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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants