From: Tobias Burnus <tobias@codesourcery.com>
Date: Tue, 24 Mar 2020 14:13:56 +0000 (+0100)
Subject: Fix OpenMP offload handling for target-link variables for nvptx (PR81689)
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=c2211a60ff05b7a0289d3e287e72c181bb4d5d8b;p=gcc.git

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