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 --git a/libgomp/target.c b/libgomp/target.c
index 6918694a843b8558dfc4fbf69ea0f803d0c8ba44..cf62af61f3b6bd4e60cd0001ca5b9e6430803c79 100644 (file)
--- 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 (file)
index 0000000..c707b38
--- /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 (file)
index 0000000..785055e
--- /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;
+    }
+}