From 9bc3b95dfefd37d860c5dc0004f8a53f6290fbb1 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Sun, 9 Feb 2020 08:17:10 +0100 Subject: [PATCH] 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 * 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. --- gcc/ChangeLog | 6 ++++++ gcc/gimplify.c | 16 ++++++++++++++++ libgomp/ChangeLog | 4 ++++ libgomp/testsuite/libgomp.c/target-38.c | 28 ++++++++++++++++++++++++++++ 4 files changed, 54 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/target-38.c 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 + + * 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 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 % 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 + + * testsuite/libgomp.c/target-38.c: New test. + 2020-02-06 Jakub Jelinek 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; +} -- 2.7.4