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 45b534b24ab41a06e3b0f24cfaaefbf5da6fd70c..ca017bedb1b4791d41501adc572df215991bcee4 100644 (file)
--- 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 b3c2138aaa31adc43bc1e0bb3705661e3fcc575b..619a42d5142330e6259fae2992449e38f9116e2f 100644 (file)
--- 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 cd34d6c9e7a95fdb450ee3ccf5a483f2b0bc97de..1c37814bde45f1a958c10ae271dabb7d7ab7ff25 100644 (file)
--- 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 11412e1059f4b16e45cc3bbe471599695132163b..c66f38b6f0cd3d16cba366a1ae5c22968fdce15c 100644 (file)
--- 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 c90cbdcc711eceb1958e31017e198aaa8d9d3271..7f5a1173eb9d46670fa14f8d862957b35977fe71 100644 (file)
--- 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 99ce33bc9b4e7eb2eeda925c137563bf42791d50..681677cc2aa90e48ca30adace57333c144140862 100644 (file)
--- 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;