| Message ID | 54a5a2ca-7b5d-a5eb-6287-ad824aebe459@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | Fix OpenMP offload handling for target-link variables for nvptx (PR81689) | expand |
On Tue, Mar 24, 2020 at 02:16:40PM +0100, Tobias Burnus wrote: > For nvptx, targetm_common.have_named_sections, hence one is You mean targetm_common.have_named_sections is false, right? > in the else branch of omp_finish_file – and the else branch > did not handle target-link variables. > > With this patch, the libgomp.c/target-link-1.c testcase now works. > > When looking at the nvptx "assembler", I saw that TREE_PUBLIC was > set for 'b' (the 'omp declare target to'-variable) but not the > target-link variables. (Not relevant for the failing test case.) > > OK for the trunk? Ok, thanks. > 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. Jakub
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/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/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;