For nvptx, targetm_common.have_named_sections, hence one is
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?
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
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;
