openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions

DECL_IN_CONSTANT_POOL are shared and thus don't really get emitted in the
BLOCK where they are used, so for OpenMP target regions that have initializers
gimplified into copying from them we actually map them at runtime from host to
offload devices.  This patch instead marks them as "omp declare target", so
that they are on the target device from the beginning and don't need to be
copied there.

2020-02-09  Jakub Jelinek  <jakub@redhat.com>

	* gimplify.c (gimplify_adjust_omp_clauses_1): Promote
	DECL_IN_CONSTANT_POOL variables into "omp declare target" to avoid
	copying them around between host and target.

	* testsuite/libgomp.c/target-38.c: New test.

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b46ea4fe14905dd4282875e8e883f39460a8343e..34c0811491a82b6cd9333b87a42fd94657ae4931 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+2020-02-09  Jakub Jelinek  <jakub@redhat.com>
+
+       * gimplify.c (gimplify_adjust_omp_clauses_1): Promote
+       DECL_IN_CONSTANT_POOL variables into "omp declare target" to avoid
+       copying them around between host and target.
+
 2020-02-08  Andrew Pinski  <apinski@marvell.com>
 
        PR target/91927

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index aafef7866373264047252a300f5d253bd493d10e..a6205d69701013e85819a336941532a6c945f3c2 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9906,6 +9906,22 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
          error ("%<_Atomic%> %qD in implicit %<map%> clause", decl);
          return 0;
        }
+      if (VAR_P (decl)
+         && DECL_IN_CONSTANT_POOL (decl)
+          && !lookup_attribute ("omp declare target",
+                               DECL_ATTRIBUTES (decl)))
+       {
+         tree id = get_identifier ("omp declare target");
+         DECL_ATTRIBUTES (decl)
+           = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+         varpool_node *node = varpool_node::get (decl);
+         if (node)
+           {
+             node->offloadable = 1;
+             if (ENABLE_OFFLOADING)
+               g->have_offload = true;
+           }
+       }
     }
   else if (flags & GOVD_SHARED)
     {

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ba14005c6ea51b5b1879e8fac7958676911318d5..0740df8b2a14d8a02f51162fa7bd41cdf16327e4 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2020-02-09  Jakub Jelinek  <jakub@redhat.com>
+
+       * testsuite/libgomp.c/target-38.c: New test.
+
 2020-02-06  Jakub Jelinek  <jakub@redhat.com>
 
        PR libgomp/93515

diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c
new file mode 100644 (file)
index 0000000..8169972
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-38.c
@@ -0,0 +1,28 @@
+#define A(n) n##0, n##1, n##2, n##3, n##4, n##5, n##6, n##7, n##8, n##9
+#define B(n) A(n##0), A(n##1), A(n##2), A(n##3), A(n##4), A(n##5), A(n##6), A(n##7), A(n##8), A(n##9)
+
+int
+foo (int x)
+{
+  int b[] = { B(4), B(5), B(6) };
+  return b[x];
+}
+
+int v[] = { 1, 2, 3, 4, 5, 6 };
+#pragma omp declare target to (foo, v)
+
+int
+main ()
+{
+  int i = 5;
+  asm ("" : "+g" (i));
+  #pragma omp target map(tofrom:i)
+  {
+    int a[] = { B(1), B(2), B(3) };
+    asm ("" : : "m" (a) : "memory");
+    i = a[i] + foo (i) + v[i & 63];
+  }
+  if (i != 105 + 405 + 6)
+    __builtin_abort ();
+  return 0;
+}