diff mbox series

[committed] openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions

Message ID 20200209072116.GF17695@tucnak
State New
Headers show
Series [committed] openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions | expand

Commit Message

Jakub Jelinek Feb. 9, 2020, 7:21 a.m. UTC 
Hi!

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.

Bootstrapped/regtested on x86_64-linux and i686-linux, tested also with
x86_64-linux to nvptx-none offloading, committed to trunk.

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.


	Jakub
diff mbox series

Patch

--- gcc/gimplify.c.jj	2020-01-15 11:05:19.248141349 +0100
+++ gcc/gimplify.c	2020-02-07 19:04:05.646917169 +0100
@@ -9902,6 +9902,22 @@  gimplify_adjust_omp_clauses_1 (splay_tre
 	  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)
     {
--- libgomp/testsuite/libgomp.c/target-38.c.jj	2020-02-07 19:06:58.717373079 +0100
+++ libgomp/testsuite/libgomp.c/target-38.c	2020-02-07 19:06:32.862753134 +0100
@@ -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;
+}