aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libgomp/ChangeLog.omp16
-rw-r--r--libgomp/libgomp.h4
-rw-r--r--libgomp/oacc-mem.c6
-rw-r--r--libgomp/target.c64
-rw-r--r--libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c74
-rw-r--r--libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c49
-rw-r--r--libgomp/testsuite/libgomp.c/target-map-zero-sized.c107
7 files changed, 304 insertions, 16 deletions
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 264a4d7..a6d676c 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -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:
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 43e7a06..a60a3d8 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -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 *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index feca3b2..a4ddb7d 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -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;
diff --git a/libgomp/target.c b/libgomp/target.c
index 796decd..ab1fae1 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -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)
{
diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c
new file mode 100644
index 0000000..3220828
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c
@@ -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);
+}
diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c
new file mode 100644
index 0000000..f968bd3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c
@@ -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);
+}
diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized.c
new file mode 100644
index 0000000..7c4ab80
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized.c
@@ -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 ();
+}
generated by cgit v1.1 at 2025-05-17 18:28:24 +0000