Mark offload symbols as global in lto
2016-02-08 Tom de Vries <tom@codesourcery.com>
PR lto/69607
* lto-partition.c (promote_offload_tables): New function.
* lto-partition.h (promote_offload_tables): Declare.
* lto.c (do_whole_program_analysis): call promote_offload_tables.
* testsuite/libgomp.c/target-36.c: New test.
* testsuite/libgomp.c/target-37.c: New test.
* testsuite/libgomp.c/target-38.c: New test.
---
gcc/lto/lto-partition.c | 28 ++++++++++
gcc/lto/lto-partition.h | 1 +
gcc/lto/lto.c | 2 +
libgomp/testsuite/libgomp.c/target-36.c | 4 ++
libgomp/testsuite/libgomp.c/target-37.c | 98 +++++++++++++++++++++++++++++++++
libgomp/testsuite/libgomp.c/target-38.c | 95 ++++++++++++++++++++++++++++++++
6 files changed, 228 insertions(+)
@@ -34,6 +34,7 @@ along with GCC; see the file COPYING3. If not see
#include "ipa-prop.h"
#include "ipa-inline.h"
#include "lto-partition.h"
+#include "omp-low.h"
vec<ltrans_partition> ltrans_partitions;
@@ -1003,6 +1004,33 @@ promote_symbol (symtab_node *node)
"Promoting as hidden: %s\n", node->name ());
}
+/* Promote the symbols in the offload tables. */
+
+void
+promote_offload_tables (void)
+{
+ if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars))
+ return;
+
+ for (unsigned i = 0; i < vec_safe_length (offload_funcs); i++)
+ {
+ tree fn_decl = (*offload_funcs)[i];
+ cgraph_node *node = cgraph_node::get (fn_decl);
+ if (node->externally_visible)
+ continue;
+ promote_symbol (node);
+ }
+
+ for (unsigned i = 0; i < vec_safe_length (offload_vars); i++)
+ {
+ tree var_decl = (*offload_vars)[i];
+ varpool_node *node = varpool_node::get (var_decl);
+ if (node->externally_visible)
+ continue;
+ promote_symbol (node);
+ }
+}
+
/* Return true if NODE needs named section even if it won't land in the partition
symbol table.
FIXME: we should really not use named sections for inline clones and master
@@ -36,6 +36,7 @@ extern vec<ltrans_partition> ltrans_partitions;
void lto_1_to_1_map (void);
void lto_max_map (void);
void lto_balanced_map (int);
+extern void promote_offload_tables (void);
void lto_promote_cross_file_statics (void);
void free_ltrans_partitions (void);
void lto_promote_statics_nonwpa (void);
@@ -3138,6 +3138,8 @@ do_whole_program_analysis (void)
to globals with hidden visibility because they are accessed from multiple
partitions. */
lto_promote_cross_file_statics ();
+ /* Promote all the offload symbols. */
+ promote_offload_tables ();
timevar_pop (TV_WHOPR_PARTITIONING);
timevar_stop (TV_PHASE_OPT_GEN);
new file mode 100644
@@ -0,0 +1,4 @@
+/* { dg-do run { target lto } } */
+/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
+
+#include "target-1.c"
new file mode 100644
@@ -0,0 +1,98 @@
+/* { dg-do run { target lto } } */
+/* { dg-additional-sources "target-38.c" } */
+/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
+
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+void
+fn1 (double *x, double *y, int z)
+{
+ int i;
+ for (i = 0; i < z; i++)
+ {
+ x[i] = i & 31;
+ y[i] = (i & 63) - 30;
+ }
+}
+
+#pragma omp declare target
+static int tgtv = 6;
+static int
+tgt (void)
+{
+ #pragma omp atomic update
+ tgtv++;
+ return 0;
+}
+#pragma omp end declare target
+
+static double
+fn2 (int x, int y, int z)
+{
+ double b[1024], c[1024], s = 0;
+ int i, j;
+ fn1 (b, c, x);
+ #pragma omp target data map(to: b)
+ {
+ #pragma omp target map(tofrom: c, s)
+ #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
+ #pragma omp distribute dist_schedule(static, 4) collapse(1)
+ for (j=0; j < x; j += y)
+ #pragma omp parallel for reduction(+:s)
+ for (i = j; i < j + y; i++)
+ tgt (), s += b[i] * c[i];
+ #pragma omp target update from(b, tgtv)
+ }
+ return s;
+}
+
+static double
+fn3 (int x)
+{
+ double b[1024], c[1024], s = 0;
+ int i;
+ fn1 (b, c, x);
+ #pragma omp target map(to: b, c) map(tofrom:s)
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ tgt (), s += b[i] * c[i];
+ return s;
+}
+
+static double
+fn4 (int x, double *p)
+{
+ double b[1024], c[1024], d[1024], s = 0;
+ int i;
+ fn1 (b, c, x);
+ fn1 (d + x, p + x, x);
+ #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \
+ map(tofrom: s)
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ s += b[i] * c[i] + d[x + i] + p[x + i];
+ return s;
+}
+
+extern int other_main (void);
+
+int
+main ()
+{
+ double a = fn2 (128, 4, 6);
+ int b = tgtv;
+ double c = fn3 (61);
+ #pragma omp target update from(tgtv)
+ int d = tgtv;
+ double e[1024];
+ double f = fn4 (64, e);
+ if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61
+ || f != 8032.0)
+ abort ();
+ other_main ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,95 @@
+/* { dg-skip-if "additional source" { *-*-* } } */
+
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+static void
+fna1 (double *x, double *y, int z)
+{
+ int i;
+ for (i = 0; i < z; i++)
+ {
+ x[i] = i & 31;
+ y[i] = (i & 63) - 30;
+ }
+}
+
+#pragma omp declare target
+static int tgtva = 6;
+static int
+tgta (void)
+{
+ #pragma omp atomic update
+ tgtva++;
+ return 0;
+}
+#pragma omp end declare target
+
+double
+fna2 (int x, int y, int z)
+{
+ double b[1024], c[1024], s = 0;
+ int i, j;
+ fna1 (b, c, x);
+ #pragma omp target data map(to: b)
+ {
+ #pragma omp target map(tofrom: c, s)
+ #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
+ #pragma omp distribute dist_schedule(static, 4) collapse(1)
+ for (j=0; j < x; j += y)
+ #pragma omp parallel for reduction(+:s)
+ for (i = j; i < j + y; i++)
+ tgta (), s += b[i] * c[i];
+ #pragma omp target update from(b, tgtva)
+ }
+ return s;
+}
+
+static double
+fna3 (int x)
+{
+ double b[1024], c[1024], s = 0;
+ int i;
+ fna1 (b, c, x);
+ #pragma omp target map(to: b, c) map(tofrom:s)
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ tgta (), s += b[i] * c[i];
+ return s;
+}
+
+static double
+fna4 (int x, double *p)
+{
+ double b[1024], c[1024], d[1024], s = 0;
+ int i;
+ fna1 (b, c, x);
+ fna1 (d + x, p + x, x);
+ #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \
+ map(tofrom: s)
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ s += b[i] * c[i] + d[x + i] + p[x + i];
+ return s;
+}
+
+extern int other_main (void);
+
+int
+other_main (void)
+{
+ double a = fna2 (128, 4, 6);
+ int b = tgtva;
+ double c = fna3 (61);
+ #pragma omp target update from(tgtva)
+ int d = tgtva;
+ double e[1024];
+ double f = fna4 (64, e);
+ if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61
+ || f != 8032.0)
+ abort ();
+ return 0;
+}