mirror of git://gcc.gnu.org/git/gcc.git
OpenMP: Fix mapping of zero-sized arrays with non-literal size: map(var[:n]), n = 0
For map(ptr[:0]), the used map kind is GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
and it is permitted that 'ptr' does not exist. 'ptr' is set to the device
pointee if it exists or to the host value otherwise.
For map(ptr[:3]), the variable is first mapped and then ptr is updated to point
to the just-mapped device data; the attachment uses GOMP_MAP_ATTACH.
For map(ptr[:n]), generates always a GOMP_MAP_ATTACH, but when n == 0, it
was failing with:
"pointer target not mapped for attach"
The solution is not to fail but first to check whether it was mapped before.
It turned out that for the mapping part, GCC adds a run-time check whether
n == 0 - and uses GOMP_MAP_ZERO_LEN_ARRAY_SECTION for the mapping.
Thus, we just have to check whether there such a mapping for the address
for which the GOMP_MAP_ATTACH. was requested. And, if there was, the
error diagnostic can be skipped.
Unsurprisingly, this issue occurs in real-world code; it was detected in
a code that distributes work via MPI and for some processes, some bounds
ended up to be zero.
libgomp/ChangeLog:
* target.c (gomp_attach_pointer): Return bool; accept additional
bool to optionally silence the fatal pointee-not-found error.
(gomp_map_vars_internal): If the pointee could not be found,
check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
* libgomp.h (gomp_attach_pointer): Update prototype.
* oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update
calls.
* testsuite/libgomp.c/target-map-zero-sized.c: New test.
* testsuite/libgomp.c/target-map-zero-sized-2.c: New test.
* testsuite/libgomp.c/target-map-zero-sized-3.c: New test.
(cherry picked from commit 814e29e390)
This commit is contained in:
Tobias Burnus
parent
9a06e4d6a1
commit
a1c4b92e57
|
|
@ -1,3 +1,19 @@
|
|||
2025-05-14 Tobias Burnus <tburnus@baylibre.com>
|
||||
|
||||
Backported from master:
|
||||
2025-05-14 Tobias Burnus <tburnus@baylibre.com>
|
||||
|
||||
* target.c (gomp_attach_pointer): Return bool; accept additional
|
||||
bool to optionally silence the fatal pointee-not-found error.
|
||||
(gomp_map_vars_internal): If the pointee could not be found,
|
||||
check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
|
||||
* libgomp.h (gomp_attach_pointer): Update prototype.
|
||||
* oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update
|
||||
calls.
|
||||
* testsuite/libgomp.c/target-map-zero-sized.c: New test.
|
||||
* testsuite/libgomp.c/target-map-zero-sized-2.c: New test.
|
||||
* testsuite/libgomp.c/target-map-zero-sized-3.c: New test.
|
||||
|
||||
2025-04-25 Thomas Schwinge <tschwinge@baylibre.com>
|
||||
|
||||
Backported from trunk:
|
||||
|
|
|
|||
|
|
@ -1487,10 +1487,10 @@ extern void gomp_copy_dev2host (struct gomp_device_descr *,
|
|||
struct goacc_asyncqueue *, void *, const void *,
|
||||
size_t);
|
||||
extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
|
||||
extern void gomp_attach_pointer (struct gomp_device_descr *,
|
||||
extern bool gomp_attach_pointer (struct gomp_device_descr *,
|
||||
struct goacc_asyncqueue *, splay_tree,
|
||||
splay_tree_key, uintptr_t, size_t,
|
||||
struct gomp_coalesce_buf *, bool);
|
||||
struct gomp_coalesce_buf *, bool, bool);
|
||||
extern void gomp_detach_pointer (struct gomp_device_descr *,
|
||||
struct goacc_asyncqueue *, splay_tree_key,
|
||||
uintptr_t, bool, struct gomp_coalesce_buf *);
|
||||
|
|
|
|||
|
|
@ -979,7 +979,7 @@ acc_attach_async (void **hostaddr, int async)
|
|||
}
|
||||
|
||||
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
|
||||
0, NULL, false);
|
||||
0, NULL, false, true);
|
||||
|
||||
gomp_mutex_unlock (&acc_dev->lock);
|
||||
}
|
||||
|
|
@ -1215,7 +1215,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
|
|||
if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
|
||||
{
|
||||
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
|
||||
(uintptr_t) h, s, NULL, false);
|
||||
(uintptr_t) h, s, NULL, false, true);
|
||||
/* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
|
||||
reference counts ('n->refcount', 'n->dynamic_refcount'). */
|
||||
}
|
||||
|
|
@ -1233,7 +1233,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
|
|||
= lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
|
||||
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
|
||||
(uintptr_t) hostaddrs[j], sizes[j], NULL,
|
||||
false);
|
||||
false, true);
|
||||
}
|
||||
|
||||
bool processed = false;
|
||||
|
|
|
|||
|
|
@ -800,12 +800,22 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
|
|||
(void *) cur_node.host_end);
|
||||
}
|
||||
|
||||
attribute_hidden void
|
||||
/* Update the devptr by setting it to the device address of the host pointee
|
||||
'attach_to'; devptr is obtained from the splay_tree_key n.
|
||||
When the pointer is already attached or the host pointee is either
|
||||
NULL or in memory map, this function returns true.
|
||||
Otherwise, the device pointer is set to point to the host pointee and:
|
||||
- If allow_zero_length_array_sections is set, true is returned.
|
||||
- Else, if fail_if_not_found is set, a fatal error is issued.
|
||||
- Otherwise, false is returned. */
|
||||
|
||||
attribute_hidden bool
|
||||
gomp_attach_pointer (struct gomp_device_descr *devicep,
|
||||
struct goacc_asyncqueue *aq, splay_tree mem_map,
|
||||
splay_tree_key n, uintptr_t attach_to, size_t bias,
|
||||
struct gomp_coalesce_buf *cbufp,
|
||||
bool allow_zero_length_array_sections)
|
||||
bool allow_zero_length_array_sections,
|
||||
bool fail_if_not_found)
|
||||
{
|
||||
struct splay_tree_key_s s;
|
||||
size_t size, idx;
|
||||
|
|
@ -860,7 +870,7 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
|
|||
gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
|
||||
sizeof (void *), true, cbufp);
|
||||
|
||||
return;
|
||||
return true;
|
||||
}
|
||||
|
||||
s.host_start = target + bias;
|
||||
|
|
@ -869,15 +879,16 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
|
|||
|
||||
if (!tn)
|
||||
{
|
||||
if (allow_zero_length_array_sections)
|
||||
/* When allowing attachment to zero-length array sections, we
|
||||
copy the host pointer when the target region is not mapped. */
|
||||
data = target;
|
||||
else
|
||||
/* We copy the host pointer when the target region is not mapped;
|
||||
for allow_zero_length_array_sections, that's permitted.
|
||||
Otherwise, it depends on the context. Return false in that
|
||||
case, unless fail_if_not_found. */
|
||||
if (!allow_zero_length_array_sections && fail_if_not_found)
|
||||
{
|
||||
gomp_mutex_unlock (&devicep->lock);
|
||||
gomp_fatal ("pointer target not mapped for attach");
|
||||
}
|
||||
data = target;
|
||||
}
|
||||
else
|
||||
data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
|
||||
|
|
@ -889,10 +900,13 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
|
|||
|
||||
gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
|
||||
sizeof (void *), true, cbufp);
|
||||
if (!tn && !allow_zero_length_array_sections)
|
||||
return false;
|
||||
}
|
||||
else
|
||||
gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
|
||||
(void *) attach_to, (int) n->aux->attach_count[idx]);
|
||||
return true;
|
||||
}
|
||||
|
||||
attribute_hidden void
|
||||
|
|
@ -1794,9 +1808,37 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
|
|||
bool zlas
|
||||
= ((kind & typemask)
|
||||
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
|
||||
gomp_attach_pointer (devicep, aq, mem_map, n,
|
||||
(uintptr_t) hostaddrs[i], sizes[i],
|
||||
cbufp, zlas);
|
||||
/* For 'target enter data', the map clauses are split;
|
||||
however, for more complex code with struct and
|
||||
pointer members, the mapping and the attach can end up
|
||||
in different sets; or the wrong mapping with the
|
||||
attach. As there is no way to know whether a size
|
||||
zero like 'var->ptr[i][:0]' happend in the same
|
||||
directive or not, the not-attached check is now
|
||||
fully silenced for 'enter data'. */
|
||||
if (openmp_p && (pragma_kind & GOMP_MAP_VARS_ENTER_DATA))
|
||||
zlas = true;
|
||||
if (!gomp_attach_pointer (devicep, aq, mem_map, n,
|
||||
(uintptr_t) hostaddrs[i], sizes[i],
|
||||
cbufp, zlas, !openmp_p))
|
||||
{
|
||||
/* Pointee not found; that's an error except for
|
||||
map(var[:n]) with n == 0; the compiler adds a
|
||||
runtime condition such that for those the kind is
|
||||
always GOMP_MAP_ZERO_LEN_ARRAY_SECTION. */
|
||||
for (j = i; j > 0; j--)
|
||||
if (*(void**) hostaddrs[i] == hostaddrs[j-1] - sizes[i]
|
||||
&& sizes[j-1] == 0
|
||||
&& (GOMP_MAP_ZERO_LEN_ARRAY_SECTION
|
||||
== (get_kind (short_mapkind, kinds, j-1)
|
||||
& typemask)))
|
||||
break;
|
||||
if (j == 0)
|
||||
{
|
||||
gomp_mutex_unlock (&devicep->lock);
|
||||
gomp_fatal ("pointer target not mapped for attach");
|
||||
}
|
||||
}
|
||||
}
|
||||
else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
|
||||
{
|
||||
|
|
|
|||
|
|
@ -0,0 +1,74 @@
|
|||
int
|
||||
main ()
|
||||
{
|
||||
int i, n;
|
||||
int data[] = {1,2};
|
||||
struct S { int **ptrset; };
|
||||
|
||||
// -----------------------------------
|
||||
|
||||
/* The produced mapping for sptr1->ptrset[i][:n]
|
||||
|
||||
GOMP_MAP_STRUCT (size = 1)
|
||||
GOMP_MAP_ZERO_LEN_ARRAY_SECTION
|
||||
GOMP_MAP_ZERO_LEN_ARRAY_SECTION
|
||||
GOMP_MAP_ATTACH
|
||||
GOMP_MAP_ATTACH -> attaching to 2nd GOMP_MAP_ZERO_LEN_ARRAY_SECTION
|
||||
|
||||
which get split into 3 separate map_vars call; in particular,
|
||||
the latter is separate and points to an unmpapped variable.
|
||||
|
||||
Thus, it failed with:
|
||||
libgomp: pointer target not mapped for attach */
|
||||
|
||||
struct S s1, *sptr1;
|
||||
s1.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3);
|
||||
s1.ptrset[0] = data;
|
||||
s1.ptrset[1] = data;
|
||||
s1.ptrset[2] = data;
|
||||
sptr1 = &s1;
|
||||
|
||||
i = 1;
|
||||
n = 0;
|
||||
#pragma omp target enter data map(sptr1[:1], sptr1->ptrset[:3])
|
||||
#pragma omp target enter data map(sptr1->ptrset[i][:n])
|
||||
|
||||
#pragma omp target exit data map(sptr1->ptrset[i][:n])
|
||||
#pragma omp target exit data map(sptr1[:1], sptr1->ptrset[:3])
|
||||
|
||||
__builtin_free (s1.ptrset);
|
||||
|
||||
// -----------------------------------
|
||||
|
||||
/* The produced mapping for sptr2->ptrset[i][:n] is similar:
|
||||
|
||||
GOMP_MAP_STRUCT (size = 1)
|
||||
GOMP_MAP_ZERO_LEN_ARRAY_SECTION
|
||||
GOMP_MAP_TO ! this one has now a finite size
|
||||
GOMP_MAP_ATTACH
|
||||
GOMP_MAP_ATTACH -> attach to the GOMP_MAP_TO
|
||||
|
||||
As the latter GOMP_MAP_ATTACH has now a pointer target,
|
||||
the attachment worked. */
|
||||
|
||||
struct S s2, *sptr2;
|
||||
s2.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3);
|
||||
s2.ptrset[0] = data;
|
||||
s2.ptrset[1] = data;
|
||||
s2.ptrset[2] = data;
|
||||
sptr2 = &s2;
|
||||
|
||||
i = 1;
|
||||
n = 2;
|
||||
#pragma omp target enter data map(sptr2[:1], sptr2->ptrset[:3])
|
||||
#pragma omp target enter data map(sptr2->ptrset[i][:n])
|
||||
|
||||
#pragma omp target
|
||||
if (sptr2->ptrset[1][0] != 1 || sptr2->ptrset[1][1] != 2)
|
||||
__builtin_abort ();
|
||||
|
||||
#pragma omp target exit data map(sptr2->ptrset[i][:n])
|
||||
#pragma omp target exit data map(sptr2[:1], sptr2->ptrset[:3])
|
||||
|
||||
__builtin_free (s2.ptrset);
|
||||
}
|
||||
|
|
@ -0,0 +1,49 @@
|
|||
int
|
||||
main ()
|
||||
{
|
||||
int i, n, n2;
|
||||
int data[] = {1,2};
|
||||
struct S {
|
||||
int **ptrset;
|
||||
int **ptrset2;
|
||||
};
|
||||
|
||||
/* This is the same as target-map-zero-sized-3.c, but by mixing
|
||||
mapped and non-mapped items, the mapping before the ATTACH
|
||||
might (or here: is) not actually associated with the the
|
||||
pointer used for attaching. Thus, if one does a simple
|
||||
|
||||
if (openmp_p
|
||||
&& (pragma_kind & GOMP_MAP_VARS_ENTER_DATA)
|
||||
&& mapnum == 1)
|
||||
check in target.c's gomp_map_vars_internal will fail
|
||||
as mapnum > 1 but still the map associated with this
|
||||
ATTACH is in a different set. */
|
||||
|
||||
struct S s1, *sptr1;
|
||||
s1.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3);
|
||||
s1.ptrset2 = (int **) __builtin_malloc (sizeof(void*) * 3);
|
||||
s1.ptrset[0] = data;
|
||||
s1.ptrset[1] = data;
|
||||
s1.ptrset[2] = data;
|
||||
s1.ptrset2[0] = data;
|
||||
s1.ptrset2[1] = data;
|
||||
s1.ptrset2[2] = data;
|
||||
sptr1 = &s1;
|
||||
|
||||
i = 1;
|
||||
n = 0;
|
||||
n2 = 2;
|
||||
#pragma omp target enter data map(sptr1[:1], sptr1->ptrset[:3], sptr1->ptrset2[:3])
|
||||
#pragma omp target enter data map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n])
|
||||
|
||||
#pragma omp target
|
||||
if (sptr1->ptrset2[1][0] != 1 || sptr1->ptrset2[1][1] != 2)
|
||||
__builtin_abort ();
|
||||
|
||||
#pragma omp target exit data map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n])
|
||||
#pragma omp target exit data map(sptr1[:1], sptr1->ptrset[:3], sptr1->ptrset2[:3])
|
||||
|
||||
__builtin_free (s1.ptrset);
|
||||
__builtin_free (s1.ptrset2);
|
||||
}
|
||||
|
|
@ -0,0 +1,107 @@
|
|||
/* { dg-do run } */
|
||||
/* { dg-additional-options "-O0" } */
|
||||
|
||||
/* Issue showed up in the real world when large data was distributed
|
||||
over multiple MPI progresses - such that for one process n == 0
|
||||
happend at run time.
|
||||
|
||||
Before map(var[:0]) and map(var[:n]) with n > 0 was handled,
|
||||
this patch now also handles map(var[:n]) with n == 0.
|
||||
|
||||
Failed before with "libgomp: pointer target not mapped for attach". */
|
||||
|
||||
/* Here, the base address is shifted - which should have no effect,
|
||||
but must work as well. */
|
||||
void
|
||||
with_offset ()
|
||||
{
|
||||
struct S {
|
||||
int *ptr1, *ptr2;
|
||||
};
|
||||
struct S s1, s2;
|
||||
int *a, *b, *c, *d;
|
||||
s1.ptr1 = (int *) 0L;
|
||||
s1.ptr2 = (int *) 0xdeedbeef;
|
||||
s2.ptr1 = (int *) 0L;
|
||||
s2.ptr2 = (int *) 0xdeedbeef;
|
||||
a = (int *) 0L;
|
||||
b = (int *) 0xdeedbeef;
|
||||
c = (int *) 0L;
|
||||
d = (int *) 0xdeedbeef;
|
||||
|
||||
int n1, n2, n3, n4;
|
||||
n1 = n2 = n3 = n4 = 0;
|
||||
|
||||
#pragma omp target enter data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
|
||||
|
||||
#pragma omp target map(s2.ptr1[4:n1], s2.ptr2[2:n2], c[6:n3], d[9:n4])
|
||||
{
|
||||
if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
|
||||
|| c != (void *) 0L || d != (void *) 0xdeedbeef)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
#pragma omp target map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
|
||||
{
|
||||
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|
||||
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|
||||
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
#pragma omp target exit data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
|
||||
}
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
struct S {
|
||||
int *ptr1, *ptr2;
|
||||
};
|
||||
struct S s1, s2;
|
||||
int *a, *b, *c, *d;
|
||||
s1.ptr1 = (int *) 0L;
|
||||
s1.ptr2 = (int *) 0xdeedbeef;
|
||||
s2.ptr1 = (int *) 0L;
|
||||
s2.ptr2 = (int *) 0xdeedbeef;
|
||||
a = (int *) 0L;
|
||||
b = (int *) 0xdeedbeef;
|
||||
c = (int *) 0L;
|
||||
d = (int *) 0xdeedbeef;
|
||||
|
||||
int n1, n2, n3, n4;
|
||||
n1 = n2 = n3 = n4 = 0;
|
||||
|
||||
#pragma omp target enter data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
|
||||
|
||||
#pragma omp target map(s2.ptr1[:n1], s2.ptr2[:n2], c[:n3], d[:n4])
|
||||
{
|
||||
if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
|
||||
|| c != (void *) 0L || d != (void *) 0xdeedbeef)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
#pragma omp target map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
|
||||
{
|
||||
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|
||||
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
|
||||
|| a != (void *) 0L || b != (void *) 0xdeedbeef)
|
||||
__builtin_abort ();
|
||||
}
|
||||
|
||||
#pragma omp target exit data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
|
||||
|
||||
with_offset ();
|
||||
}
|
||||
Loading…
Reference in New Issue