diff options
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/gimplify.c | 16 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 4 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/target-38.c | 28 |
4 files changed, 54 insertions, 0 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index b46ea4f..34c0811 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 aafef78..a6205d6 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 ba14005..0740df8 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 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; +} |