https://gcc.gnu.org/g:4cb20dc043cf70b8a1b4846c86599cc1ff9680d9
commit r15-3836-g4cb20dc043cf70b8a1b4846c86599cc1ff9680d9
Author: Tobias Burnus <tbur...@baylibre.com>
Date: Tue Sep 24 17:41:39 2024 +0200
libgomp: with USM, init 'link' variables with host address
If requires unified_shared_memory or self_maps is set, make
'declare target link' variables to point initially to the host pointer.
libgomp/ChangeLog:
* target.c (gomp_load_image_to_device): For requires
unified_shared_memory, update 'link' vars to point to the host var.
* testsuite/libgomp.c-c++-common/target-link-3.c: New test.
* testsuite/libgomp.c-c++-common/target-link-4.c: New test.
Diff:
---
libgomp/target.c | 6 +++
.../testsuite/libgomp.c-c++-common/target-link-3.c | 52 ++++++++++++++++++++++
.../testsuite/libgomp.c-c++-common/target-link-4.c | 52 ++++++++++++++++++++++
3 files changed, 110 insertions(+)
diff --git a/libgomp/target.c b/libgomp/target.c
index 6918694a843b..cf62af61f3b6 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2454,6 +2454,12 @@ gomp_load_image_to_device (struct gomp_device_descr
*devicep, unsigned version,
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
array++;
+
+ if (is_link_var
+ && (omp_requires_mask
+ & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY |
GOMP_REQUIRES_SELF_MAPS)))
+ gomp_copy_host2dev (devicep, NULL, (void *) target_var->start,
+ &k->host_start, sizeof (void *), false, NULL);
}
/* Last entry is for the ICV struct variable; if absent, start = end = 0. */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c
b/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c
new file mode 100644
index 000000000000..c707b38b7d46
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-3.c
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <stdint.h>
+#include <omp.h>
+
+#pragma omp requires unified_shared_memory
+
+int A[3] = {-3,-4,-5};
+static int q = -401;
+#pragma omp declare target link(A, q)
+
+#pragma omp begin declare target
+void
+f (uintptr_t *pA, uintptr_t *pq)
+{
+ if (A[0] != 1 || A[1] != 2 || A[2] != 3 || q != 42)
+ __builtin_abort ();
+ A[0] = 13;
+ A[1] = 14;
+ A[2] = 15;
+ q = 23;
+ *pA = (uintptr_t) &A[0];
+ *pq = (uintptr_t) &q;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+ uintptr_t hpA = (uintptr_t) &A[0];
+ uintptr_t hpq = (uintptr_t) &q;
+ uintptr_t dpA, dpq;
+
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+
+ for (int i = 0; i <= omp_get_num_devices (); ++i)
+ {
+ #pragma omp target device(device_num: i) map(dpA, dpq)
+ f (&dpA, &dpq);
+ if (hpA != dpA || hpq != dpq)
+ __builtin_abort ();
+ if (A[0] != 13 || A[1] != 14 || A[2] != 15 || q != 23)
+ __builtin_abort ();
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+ }
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c
b/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c
new file mode 100644
index 000000000000..785055e216d7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-4.c
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <stdint.h>
+#include <omp.h>
+
+#pragma omp requires self_maps
+
+int A[3] = {-3,-4,-5};
+static int q = -401;
+#pragma omp declare target link(A, q)
+
+#pragma omp begin declare target
+void
+f (uintptr_t *pA, uintptr_t *pq)
+{
+ if (A[0] != 1 || A[1] != 2 || A[2] != 3 || q != 42)
+ __builtin_abort ();
+ A[0] = 13;
+ A[1] = 14;
+ A[2] = 15;
+ q = 23;
+ *pA = (uintptr_t) &A[0];
+ *pq = (uintptr_t) &q;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+ uintptr_t hpA = (uintptr_t) &A[0];
+ uintptr_t hpq = (uintptr_t) &q;
+ uintptr_t dpA, dpq;
+
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+
+ for (int i = 0; i <= omp_get_num_devices (); ++i)
+ {
+ #pragma omp target device(device_num: i) map(dpA, dpq)
+ f (&dpA, &dpq);
+ if (hpA != dpA || hpq != dpq)
+ __builtin_abort ();
+ if (A[0] != 13 || A[1] != 14 || A[2] != 15 || q != 23)
+ __builtin_abort ();
+ A[0] = 1;
+ A[1] = 2;
+ A[2] = 3;
+ q = 42;
+ }
+}
