[PATCH, PR69607] Mark offload symbols as global in lto

[PATCH, PR69607] Mark offload symbols as global in lto

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 734 bytes --]

Hi,

when running libgomp.c testsuite with "-flto -flto-partition=1to1
-fno-toplevel-reorder" we run into many compilation failures like this:
...
/tmp/xxxxxxxx.ltrans0.ltrans.o:(.gnu.offload_funcs+0x1a0): undefined 
reference to `MAIN__._omp_fn.0'^M
...

The problem is that the offload table is in one lto partition, and the 
function listed in the offload table is in another, without the function 
having been promoted to be visible in the other partition.

The patch fixes this by promoting the symbols in the offload table such 
that they're visible in all partitions.

Bootstrapped and reg-tested on x86_64.

Build for nvidia accelerator and reg-tested libgomp with various lto 
settings.

OK for trunk, stage1?

Thanks,
- Tom

[-- Attachment #2: 0005-Mark-offload-symbols-as-global-in-lto.patch --]
[-- Type: text/x-patch, Size: 3256 bytes --]

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.

---
 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 ++++
 4 files changed, 35 insertions(+)

diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
index 9eb63c2..56598d4 100644
--- a/gcc/lto/lto-partition.c
+++ b/gcc/lto/lto-partition.c
@@ -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
diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h
index 31e3764..1a38126 100644
--- a/gcc/lto/lto-partition.h
+++ b/gcc/lto/lto-partition.h
@@ -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);
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 6718fbbe..ecec1bc 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c
new file mode 100644
index 0000000..bafb718
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-36.c
@@ -0,0 +1,4 @@
+/* { dg-do run { target lto } } */
+/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
+
+#include "target-1.c"

Re: [PATCH, PR69607] Mark offload symbols as global in lto

[Ilya Verbin]

On Mon, Feb 08, 2016 at 14:00:00 +0100, Tom de Vries wrote:
> when running libgomp.c testsuite with "-flto -flto-partition=1to1
> -fno-toplevel-reorder" we run into many compilation failures like this:
> ...
> /tmp/xxxxxxxx.ltrans0.ltrans.o:(.gnu.offload_funcs+0x1a0): undefined
> reference to `MAIN__._omp_fn.0'^M
> ...
> 
> The problem is that the offload table is in one lto partition, and the
> function listed in the offload table is in another, without the function
> having been promoted to be visible in the other partition.
> 
> The patch fixes this by promoting the symbols in the offload table such that
> they're visible in all partitions.
> 
> Bootstrapped and reg-tested on x86_64.
> 
> Build for nvidia accelerator and reg-tested libgomp with various lto
> settings.

Works fine with intelmic offloading.

  -- Ilya

Re: [PATCH, PR69607] Mark offload symbols as global in lto

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 938 bytes --]

On 08/02/16 14:00, Tom de Vries wrote:
> Hi,
>
> when running libgomp.c testsuite with "-flto -flto-partition=1to1
> -fno-toplevel-reorder" we run into many compilation failures like this:
> ...
> /tmp/xxxxxxxx.ltrans0.ltrans.o:(.gnu.offload_funcs+0x1a0): undefined
> reference to `MAIN__._omp_fn.0'^M
> ...
>
> The problem is that the offload table is in one lto partition, and the
> function listed in the offload table is in another, without the function
> having been promoted to be visible in the other partition.
>
> The patch fixes this by promoting the symbols in the offload table such
> that they're visible in all partitions.
>
> Bootstrapped and reg-tested on x86_64.
>
> Build for nvidia accelerator and reg-tested libgomp with various lto
> settings.
>

Added multi-source testcase target-3{7,8}.c that triggers the PR for 
intelmicemul accelerator.

OK for trunk, stage1 (or stage4, if that's appropriate)?

Thanks,
- Tom


[-- Attachment #2: 0001-Mark-offload-symbols-as-global-in-lto.patch --]
[-- Type: text/x-patch, Size: 8045 bytes --]

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(+)

diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
index 9eb63c2..56598d4 100644
--- a/gcc/lto/lto-partition.c
+++ b/gcc/lto/lto-partition.c
@@ -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
diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h
index 31e3764..1a38126 100644
--- a/gcc/lto/lto-partition.h
+++ b/gcc/lto/lto-partition.h
@@ -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);
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 9dd513f..2736c5c 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c
new file mode 100644
index 0000000..bafb718
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-36.c
@@ -0,0 +1,4 @@
+/* { dg-do run { target lto } } */
+/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
+
+#include "target-1.c"
diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
new file mode 100644
index 0000000..1edb21e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-37.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c
new file mode 100644
index 0000000..15e69c4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-38.c
@@ -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;
+}

Re: [PATCH, PR69607] Mark offload symbols as global in lto

[Jakub Jelinek]

On Wed, Feb 17, 2016 at 01:02:17PM +0100, Tom de Vries wrote:
> Mark offload symbols as global in lto

I'm really not familiar with that part of LTO, so I'm CCing Honza and
Richard here.

> 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.

Just one space instead of two after :

> 	* lto.c (do_whole_program_analysis): call promote_offload_tables.

Capital C in Call.

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-36.c
> @@ -0,0 +1,4 @@
> +/* { dg-do run { target lto } } */
> +/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
> +
> +#include "target-1.c"
> diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
> new file mode 100644
> index 0000000..1edb21e
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-37.c
> @@ -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);

Why the C++ stuff in there?  Do you intend to include the testcase
also in libgomp.c++?  If not, it is not needed.
Otherwise, the tests LGTM.

	Jakub

Re: [PATCH, PR69607] Mark offload symbols as global in lto

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1369 bytes --]

On 17/02/16 13:30, Jakub Jelinek wrote:
> On Wed, Feb 17, 2016 at 01:02:17PM +0100, Tom de Vries wrote:
>> Mark offload symbols as global in lto
>
> I'm really not familiar with that part of LTO, so I'm CCing Honza and
> Richard here.
>> 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.
>
> Just one space instead of two after :
>
>> 	* lto.c (do_whole_program_analysis): call promote_offload_tables.
>
> Capital C in Call.
>

Done.

>> diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
>> new file mode 100644
>> index 0000000..1edb21e
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/target-37.c
>> @@ -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);
>
> Why the C++ stuff in there?  Do you intend to include the testcase
> also in libgomp.c++?

No, that's just there because I started both target-37.c and target-38.c 
by copying target-1.c.

> If not, it is not needed.

Removed.

> Otherwise, the tests LGTM.
>

Updated patch attached.

Thanks,
- Tom


[-- Attachment #2: 0001-Mark-offload-symbols-as-global-in-lto.patch --]
[-- Type: text/x-patch, Size: 7975 bytes --]

Mark offload symbols as global in lto

2016-02-17  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 | 94 +++++++++++++++++++++++++++++++++
 libgomp/testsuite/libgomp.c/target-38.c | 91 +++++++++++++++++++++++++++++++
 6 files changed, 220 insertions(+)

diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
index 9eb63c2..56598d4 100644
--- a/gcc/lto/lto-partition.c
+++ b/gcc/lto/lto-partition.c
@@ -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
diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h
index 31e3764..1a38126 100644
--- a/gcc/lto/lto-partition.h
+++ b/gcc/lto/lto-partition.h
@@ -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);
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 9dd513f..2736c5c 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -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);
diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c
new file mode 100644
index 0000000..bafb718
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-36.c
@@ -0,0 +1,4 @@
+/* { dg-do run { target lto } } */
+/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
+
+#include "target-1.c"
diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
new file mode 100644
index 0000000..fe5b8ef
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-37.c
@@ -0,0 +1,94 @@
+/* { dg-do run { target lto } } */
+/* { dg-additional-sources "target-38.c" } */
+/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
+
+extern 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;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c
new file mode 100644
index 0000000..de2b4c8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-38.c
@@ -0,0 +1,91 @@
+/* { dg-skip-if "additional source" { *-*-* } } */
+
+extern 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;
+}

[PING][PATCH, PR69607] Mark offload symbols as global in lto

[Tom de Vries]

On 17/02/16 16:48, Tom de Vries wrote:
> On 17/02/16 13:30, Jakub Jelinek wrote:
>> On Wed, Feb 17, 2016 at 01:02:17PM +0100, Tom de Vries wrote:
>>> Mark offload symbols as global in lto
>>
>> I'm really not familiar with that part of LTO, so I'm CCing Honza and
>> Richard here.
>>> 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.
>>
>> Just one space instead of two after :
>>
>>>     * lto.c (do_whole_program_analysis): call promote_offload_tables.
>>
>> Capital C in Call.
>>
>
> Done.
>
>>> diff --git a/libgomp/testsuite/libgomp.c/target-37.c
>>> b/libgomp/testsuite/libgomp.c/target-37.c
>>> new file mode 100644
>>> index 0000000..1edb21e
>>> --- /dev/null
>>> +++ b/libgomp/testsuite/libgomp.c/target-37.c
>>> @@ -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);
>>
>> Why the C++ stuff in there?  Do you intend to include the testcase
>> also in libgomp.c++?
>
> No, that's just there because I started both target-37.c and target-38.c
> by copying target-1.c.
>
>> If not, it is not needed.
>
> Removed.
>
>> Otherwise, the tests LGTM.
>>
>
> Updated patch attached.
>

Ping.

[ Original submission here:
     https://gcc.gnu.org/ml/gcc-patches/2016-02/msg00547.html
   Latest patch with updated testcases:
     https://gcc.gnu.org/ml/gcc-patches/2016-02/msg01196.html
]

OK for trunk, stage1 (or stage4, if that's appropriate)?

> Thanks,
> - Tom
>
>
> 0001-Mark-offload-symbols-as-global-in-lto.patch
>
>
> Mark offload symbols as global in lto
>
> 2016-02-17  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 | 94 +++++++++++++++++++++++++++++++++
>   libgomp/testsuite/libgomp.c/target-38.c | 91 +++++++++++++++++++++++++++++++
>   6 files changed, 220 insertions(+)
>
> diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
> index 9eb63c2..56598d4 100644
> --- a/gcc/lto/lto-partition.c
> +++ b/gcc/lto/lto-partition.c
> @@ -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
> diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h
> index 31e3764..1a38126 100644
> --- a/gcc/lto/lto-partition.h
> +++ b/gcc/lto/lto-partition.h
> @@ -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);
> diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
> index 9dd513f..2736c5c 100644
> --- a/gcc/lto/lto.c
> +++ b/gcc/lto/lto.c
> @@ -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);
> diff --git a/libgomp/testsuite/libgomp.c/target-36.c b/libgomp/testsuite/libgomp.c/target-36.c
> new file mode 100644
> index 0000000..bafb718
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-36.c
> @@ -0,0 +1,4 @@
> +/* { dg-do run { target lto } } */
> +/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
> +
> +#include "target-1.c"
> diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
> new file mode 100644
> index 0000000..fe5b8ef
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-37.c
> @@ -0,0 +1,94 @@
> +/* { dg-do run { target lto } } */
> +/* { dg-additional-sources "target-38.c" } */
> +/* { dg-additional-options "-flto -flto-partition=1to1 -fno-toplevel-reorder" } */
> +
> +extern 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;
> +}
> diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c
> new file mode 100644
> index 0000000..de2b4c8
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-38.c
> @@ -0,0 +1,91 @@
> +/* { dg-skip-if "additional source" { *-*-* } } */
> +
> +extern 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;
> +}
>

[PING^2][PATCH, PR69607] Mark offload symbols as global in lto

[Tom de Vries]

On 24/02/16 14:37, Tom de Vries wrote:
> On 17/02/16 16:48, Tom de Vries wrote:
>> On 17/02/16 13:30, Jakub Jelinek wrote:
>>> On Wed, Feb 17, 2016 at 01:02:17PM +0100, Tom de Vries wrote:
>>>> Mark offload symbols as global in lto
>>>
>>> I'm really not familiar with that part of LTO, so I'm CCing Honza and
>>> Richard here.
>>>> 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.
>>>
>>> Just one space instead of two after :
>>>
>>>>     * lto.c (do_whole_program_analysis): call promote_offload_tables.
>>>
>>> Capital C in Call.
>>>
>>
>> Done.
>>
>>>> diff --git a/libgomp/testsuite/libgomp.c/target-37.c
>>>> b/libgomp/testsuite/libgomp.c/target-37.c
>>>> new file mode 100644
>>>> index 0000000..1edb21e
>>>> --- /dev/null
>>>> +++ b/libgomp/testsuite/libgomp.c/target-37.c
>>>> @@ -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);
>>>
>>> Why the C++ stuff in there?  Do you intend to include the testcase
>>> also in libgomp.c++?
>>
>> No, that's just there because I started both target-37.c and target-38.c
>> by copying target-1.c.
>>
>>> If not, it is not needed.
>>
>> Removed.
>>
>>> Otherwise, the tests LGTM.
>>>
>>
>> Updated patch attached.
>>
>

Ping^2.

This patch fixes an ICE when compiling an openmp program using -flto 
-partition=1to1 -fno-toplevel-reorder.

[ Original submission here:
      https://gcc.gnu.org/ml/gcc-patches/2016-02/msg00547.html
    Latest patch with updated testcases:
      https://gcc.gnu.org/ml/gcc-patches/2016-02/msg01196.html
]

OK for trunk, stage1 (or stage4, if that's appropriate)?

Thanks,
- Tom

>>
>>
>> 0001-Mark-offload-symbols-as-global-in-lto.patch
>>
>>
>> Mark offload symbols as global in lto
>>
>> 2016-02-17  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 | 94
>> +++++++++++++++++++++++++++++++++
>>   libgomp/testsuite/libgomp.c/target-38.c | 91
>> +++++++++++++++++++++++++++++++
>>   6 files changed, 220 insertions(+)
>>
>> diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
>> index 9eb63c2..56598d4 100644
>> --- a/gcc/lto/lto-partition.c
>> +++ b/gcc/lto/lto-partition.c
>> @@ -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
>> diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h
>> index 31e3764..1a38126 100644
>> --- a/gcc/lto/lto-partition.h
>> +++ b/gcc/lto/lto-partition.h
>> @@ -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);
>> diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
>> index 9dd513f..2736c5c 100644
>> --- a/gcc/lto/lto.c
>> +++ b/gcc/lto/lto.c
>> @@ -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);
>> diff --git a/libgomp/testsuite/libgomp.c/target-36.c
>> b/libgomp/testsuite/libgomp.c/target-36.c
>> new file mode 100644
>> index 0000000..bafb718
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/target-36.c
>> @@ -0,0 +1,4 @@
>> +/* { dg-do run { target lto } } */
>> +/* { dg-additional-options "-flto -flto-partition=1to1
>> -fno-toplevel-reorder" } */
>> +
>> +#include "target-1.c"
>> diff --git a/libgomp/testsuite/libgomp.c/target-37.c
>> b/libgomp/testsuite/libgomp.c/target-37.c
>> new file mode 100644
>> index 0000000..fe5b8ef
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/target-37.c
>> @@ -0,0 +1,94 @@
>> +/* { dg-do run { target lto } } */
>> +/* { dg-additional-sources "target-38.c" } */
>> +/* { dg-additional-options "-flto -flto-partition=1to1
>> -fno-toplevel-reorder" } */
>> +
>> +extern 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;
>> +}
>> diff --git a/libgomp/testsuite/libgomp.c/target-38.c
>> b/libgomp/testsuite/libgomp.c/target-38.c
>> new file mode 100644
>> index 0000000..de2b4c8
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/target-38.c
>> @@ -0,0 +1,91 @@
>> +/* { dg-skip-if "additional source" { *-*-* } } */
>> +
>> +extern 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;
>> +}
>>
>

[PING^3][PATCH, PR69607] Mark offload symbols as global in lto

[Tom de Vries]

On 07/03/16 15:38, Tom de Vries wrote:
> On 24/02/16 14:37, Tom de Vries wrote:
>> On 17/02/16 16:48, Tom de Vries wrote:
>>> On 17/02/16 13:30, Jakub Jelinek wrote:
>>>> On Wed, Feb 17, 2016 at 01:02:17PM +0100, Tom de Vries wrote:
>>>>> Mark offload symbols as global in lto
>>>>
>>>> I'm really not familiar with that part of LTO, so I'm CCing Honza and
>>>> Richard here.
>>>>> 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.
>>>>
>>>> Just one space instead of two after :
>>>>
>>>>>     * lto.c (do_whole_program_analysis): call promote_offload_tables.
>>>>
>>>> Capital C in Call.
>>>>
>>>
>>> Done.
>>>
>>>>> diff --git a/libgomp/testsuite/libgomp.c/target-37.c
>>>>> b/libgomp/testsuite/libgomp.c/target-37.c
>>>>> new file mode 100644
>>>>> index 0000000..1edb21e
>>>>> --- /dev/null
>>>>> +++ b/libgomp/testsuite/libgomp.c/target-37.c
>>>>> @@ -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);
>>>>
>>>> Why the C++ stuff in there?  Do you intend to include the testcase
>>>> also in libgomp.c++?
>>>
>>> No, that's just there because I started both target-37.c and target-38.c
>>> by copying target-1.c.
>>>
>>>> If not, it is not needed.
>>>
>>> Removed.
>>>
>>>> Otherwise, the tests LGTM.
>>>>
>>>
>>> Updated patch attached.
>>>
>>

Ping^3.


This patch fixes an ICE when compiling an openmp program using -flto
-partition=1to1 -fno-toplevel-reorder.

[ Original submission here:
       https://gcc.gnu.org/ml/gcc-patches/2016-02/msg00547.html
     Latest patch with updated testcases:
       https://gcc.gnu.org/ml/gcc-patches/2016-02/msg01196.html
]

OK for trunk, stage1 (or stage4, if that's appropriate)?

Thanks,
- Tom

>>>
>>>
>>> 0001-Mark-offload-symbols-as-global-in-lto.patch
>>>
>>>
>>> Mark offload symbols as global in lto
>>>
>>> 2016-02-17  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 | 94
>>> +++++++++++++++++++++++++++++++++
>>>   libgomp/testsuite/libgomp.c/target-38.c | 91
>>> +++++++++++++++++++++++++++++++
>>>   6 files changed, 220 insertions(+)
>>>
>>> diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
>>> index 9eb63c2..56598d4 100644
>>> --- a/gcc/lto/lto-partition.c
>>> +++ b/gcc/lto/lto-partition.c
>>> @@ -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
>>> diff --git a/gcc/lto/lto-partition.h b/gcc/lto/lto-partition.h
>>> index 31e3764..1a38126 100644
>>> --- a/gcc/lto/lto-partition.h
>>> +++ b/gcc/lto/lto-partition.h
>>> @@ -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);
>>> diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
>>> index 9dd513f..2736c5c 100644
>>> --- a/gcc/lto/lto.c
>>> +++ b/gcc/lto/lto.c
>>> @@ -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);
>>> diff --git a/libgomp/testsuite/libgomp.c/target-36.c
>>> b/libgomp/testsuite/libgomp.c/target-36.c
>>> new file mode 100644
>>> index 0000000..bafb718
>>> --- /dev/null
>>> +++ b/libgomp/testsuite/libgomp.c/target-36.c
>>> @@ -0,0 +1,4 @@
>>> +/* { dg-do run { target lto } } */
>>> +/* { dg-additional-options "-flto -flto-partition=1to1
>>> -fno-toplevel-reorder" } */
>>> +
>>> +#include "target-1.c"
>>> diff --git a/libgomp/testsuite/libgomp.c/target-37.c
>>> b/libgomp/testsuite/libgomp.c/target-37.c
>>> new file mode 100644
>>> index 0000000..fe5b8ef
>>> --- /dev/null
>>> +++ b/libgomp/testsuite/libgomp.c/target-37.c
>>> @@ -0,0 +1,94 @@
>>> +/* { dg-do run { target lto } } */
>>> +/* { dg-additional-sources "target-38.c" } */
>>> +/* { dg-additional-options "-flto -flto-partition=1to1
>>> -fno-toplevel-reorder" } */
>>> +
>>> +extern 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;
>>> +}
>>> diff --git a/libgomp/testsuite/libgomp.c/target-38.c
>>> b/libgomp/testsuite/libgomp.c/target-38.c
>>> new file mode 100644
>>> index 0000000..de2b4c8
>>> --- /dev/null
>>> +++ b/libgomp/testsuite/libgomp.c/target-38.c
>>> @@ -0,0 +1,91 @@
>>> +/* { dg-skip-if "additional source" { *-*-* } } */
>>> +
>>> +extern 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;
>>> +}
>>>
>>
>