From: Jakub Jelinek <jakub@redhat.com>
Date: Sun, 9 Feb 2020 07:17:10 +0000 (+0100)
Subject: openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions
X-Git-Tag: basepoints/gcc-11~1546
X-Git-Url: http://git.ipfire.org/?a=commitdiff_plain;h=9bc3b95dfefd37d860c5dc0004f8a53f6290fbb1;p=thirdparty%2Fgcc.git

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 b46ea4fe1490..34c0811491a8 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 aafef7866373..a6205d697010 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 ba14005c6ea5..0740df8b2a14 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 000000000000..816997205267
--- /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;
+}