b62b58d1bb
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.
30 lines
549 B
C++
30 lines
549 B
C++
// 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;
|
|
}
|