[all-commits] [llvm/llvm-project] b62b58: [OpenMP] Fix crash with duplicate mapping on targe...

[Julian Brown via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: b62b58d1bbbff7ca200f166603c80470639a0632
      https://github.com/llvm/llvm-project/commit/b62b58d1bbbff7ca200f166603c80470639a0632
  Author: Julian Brown <julian.brown at amd.com>
  Date:   2025-06-29 (Sun, 29 Jun 2025)

  Changed paths:
    M offload/libomptarget/OpenMP/Mapping.cpp
    A offload/test/mapping/duplicate_mappings_1.cpp
    A offload/test/mapping/duplicate_mappings_2.cpp

  Log Message:
  -----------
  [OpenMP] Fix crash with duplicate mapping on target directive (#146136)

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.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications