[gcc(refs/vendors/redhat/heads/gcc-10-branch)] Fix OpenMP offload handling for target-link variables for nvptx (PR81689)

[Jakub Jelinek <jakub@gcc.gnu.org>]

https://gcc.gnu.org/g:c2211a60ff05b7a0289d3e287e72c181bb4d5d8b

commit c2211a60ff05b7a0289d3e287e72c181bb4d5d8b
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Tue Mar 24 15:13:56 2020 +0100

    Fix OpenMP offload handling for target-link variables for nvptx (PR81689)
    
            PR libgomp/81689
            * lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state.
    
            PR libgomp/81689
            * omp-offload.c (omp_finish_file): Fix target-link handling if
            targetm_common.have_named_sections is false.
    
            PR libgomp/81689
            * testsuite/libgomp.c/target-link-1.c: Remove xfail.

Diff:
---
 gcc/ChangeLog                               |  6 ++++++
 gcc/lto/ChangeLog                           |  5 +++++
 gcc/lto/lto.c                               |  1 +
 gcc/omp-offload.c                           | 14 +++++++++++++-
 libgomp/ChangeLog                           |  5 +++++
 libgomp/testsuite/libgomp.c/target-link-1.c |  3 ---
 6 files changed, 30 insertions(+), 4 deletions(-)

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 45b534b24ab..ca017bedb1b 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+2020-03-24  Tobias Burnus  <tobias@codesourcery.com>
+
+	PR libgomp/81689
+	* omp-offload.c (omp_finish_file): Fix target-link handling if
+	targetm_common.have_named_sections is false.
+
 2020-03-24  Jakub Jelinek  <jakub@redhat.com>
 
 	PR target/94286
diff --git a/gcc/lto/ChangeLog b/gcc/lto/ChangeLog
index b3c2138aaa3..619a42d5142 100644
--- a/gcc/lto/ChangeLog
+++ b/gcc/lto/ChangeLog
@@ -1,3 +1,8 @@
+2020-03-24  Tobias Burnus  <tobias@codesourcery.com>
+
+	PR libgomp/81689
+	* lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state.
+
 2020-01-29  Tobias Burnus  <tobias@codesourcery.com>
 
 	* lto.c (offload_handle_link_vars): Reduce chance of
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index cd34d6c9e7a..1c37814bde4 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -566,6 +566,7 @@ offload_handle_link_vars (void)
 							     "linkptr"), type);
 	TREE_USED (link_ptr_var) = 1;
 	TREE_STATIC (link_ptr_var) = 1;
+	TREE_PUBLIC (link_ptr_var) = TREE_PUBLIC (var->decl);
 	DECL_ARTIFICIAL (link_ptr_var) = 1;
 	SET_DECL_ASSEMBLER_NAME (link_ptr_var, DECL_NAME (link_ptr_var));
 	SET_DECL_VALUE_EXPR (var->decl, build_simple_mem_ref (link_ptr_var));
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 11412e1059f..c66f38b6f0c 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -220,7 +220,19 @@ omp_finish_file (void)
       for (unsigned i = 0; i < num_vars; i++)
 	{
 	  tree it = (*offload_vars)[i];
-	  targetm.record_offload_symbol (it);
+#ifdef ACCEL_COMPILER
+	  if (DECL_HAS_VALUE_EXPR_P (it)
+	      && lookup_attribute ("omp declare target link",
+				   DECL_ATTRIBUTES (it)))
+	    {
+	      tree value_expr = DECL_VALUE_EXPR (it);
+	      tree link_ptr_decl = TREE_OPERAND (value_expr, 0);
+	      targetm.record_offload_symbol (link_ptr_decl);
+	      varpool_node::finalize_decl (link_ptr_decl);
+	    }
+	  else
+#endif
+	    targetm.record_offload_symbol (it);
 	}
     }
 }
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index c90cbdcc711..7f5a1173eb9 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2020-03-24  Tobias Burnus  <tobias@codesourcery.com>
+
+	PR libgomp/81689
+	* testsuite/libgomp.c/target-link-1.c: Remove xfail.
+
 2020-03-20  Tobias Burnus  <tobias@codesourcery.com>
 
 	PR libgomp/94251
diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c b/libgomp/testsuite/libgomp.c/target-link-1.c
index 99ce33bc9b4..681677cc2aa 100644
--- a/libgomp/testsuite/libgomp.c/target-link-1.c
+++ b/libgomp/testsuite/libgomp.c/target-link-1.c
@@ -1,6 +1,3 @@
-/* { dg-xfail-if "#pragma omp target link not implemented" { offload_target_nvptx } }
-   Cf. https://gcc.gnu.org/PR81689.  */
-
 struct S { int s, t; };
 
 int a = 1, b = 1;