From c2211a60ff05b7a0289d3e287e72c181bb4d5d8b Mon Sep 17 00:00:00 2001 From: Tobias Burnus Date: Tue, 24 Mar 2020 15:13:56 +0100 Subject: [PATCH] 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. --- 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 + + 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 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 + + PR libgomp/81689 + * lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state. + 2020-01-29 Tobias Burnus * 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 + + PR libgomp/81689 + * testsuite/libgomp.c/target-link-1.c: Remove xfail. + 2020-03-20 Tobias Burnus 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; -- 2.30.2