Commit 9bc3b95d by Jakub Jelinek

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.
parent a5691173
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> 2020-02-08 Andrew Pinski <apinski@marvell.com>
PR target/91927 PR target/91927
......
...@@ -9906,6 +9906,22 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) ...@@ -9906,6 +9906,22 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
error ("%<_Atomic%> %qD in implicit %<map%> clause", decl); error ("%<_Atomic%> %qD in implicit %<map%> clause", decl);
return 0; 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) else if (flags & GOVD_SHARED)
{ {
......
2020-02-09 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c/target-38.c: New test.
2020-02-06 Jakub Jelinek <jakub@redhat.com> 2020-02-06 Jakub Jelinek <jakub@redhat.com>
PR libgomp/93515 PR libgomp/93515
......
#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;
}
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment