mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-02-27 12:05:56 +08:00
Code Issues Packages Projects Releases Wiki Activity

Fix OpenMP offload handling for target-link variables for nvptx (PR81689)

Browse Source 
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.
This commit is contained in:
Tobias Burnus 2020-03-24 15:13:56 +01:00
parent 906b3eb9df
commit c2211a60ff
6 changed files with 30 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
gcc/ChangeLog
View File

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

5
gcc/lto/ChangeLog
View File

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

1
gcc/lto/lto.c
View File

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

14
gcc/omp-offload.c
View File

@ -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);
}
}
}

5
libgomp/ChangeLog
View File

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

3
libgomp/testsuite/libgomp.c/target-link-1.c
View File

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