[gcc(refs/users/marxin/heads/marxin-gcc-benchmark-branch)] openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions

[Martin Liska]

https://gcc.gnu.org/g:9bc3b95dfefd37d860c5dc0004f8a53f6290fbb1

commit 9bc3b95dfefd37d860c5dc0004f8a53f6290fbb1
Author: Jakub Jelinek <jakub@redhat.com>
Date:   Sun Feb 9 08:17:10 2020 +0100

    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:
---
 gcc/ChangeLog                           |  6 ++++++
 gcc/gimplify.c                          | 16 ++++++++++++++++
 libgomp/ChangeLog                       |  4 ++++
 libgomp/testsuite/libgomp.c/target-38.c | 28 ++++++++++++++++++++++++++++
 4 files changed, 54 insertions(+)

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b46ea4fe149..34c0811491a 100644
--- 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 aafef786637..a6205d69701 100644
--- 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 ba14005c6ea..0740df8b2a1 100644
--- 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
index 00000000000..81699720526
--- /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;
+}