LTO: Fixes for renaming issues with offload/OpenMP [PR104285]

gcc/lto/ChangeLog:

	PR middle-end/104285
	* lto-partition.c (maybe_rewrite_identifier): Use get_identifier
	for the returned string to be usable as hash key.
	(validize_symbol_for_target): Hence, use return value directly.
	(privatize_symbol_name_1): Track maybe_rewrite_identifier renames.
	* lto.c (offload_handle_link_vars): Move function up before ...
	(do_whole_program_analysis): Call it after static renamings.
	(lto_main): Move call after static renamings.

libgomp/ChangeLog:

	PR middle-end/104285
	* testsuite/libgomp.c++/target-same-name-2-a.C: New test.
	* testsuite/libgomp.c++/target-same-name-2-b.C: New test.
	* testsuite/libgomp.c++/target-same-name-2.C: New test.
	* testsuite/libgomp.c-c++-common/target-same-name-1-a.c: New test.
	* testsuite/libgomp.c-c++-common/target-same-name-1-b.c: New test.
	* testsuite/libgomp.c-c++-common/target-same-name-1.c: New test.

(cherry picked from commit 1002a7ace111d746249fdea71af9b8e039cea0eb)

8 files changed, 331 insertions, 34 deletions

diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
index 15761ac9eb5..67dc0a595fd 100644
--- a/gcc/lto/lto-partition.c
+++ b/gcc/lto/lto-partition.c
@@ -896,6 +896,11 @@ maybe_rewrite_identifier (const char *ptr)
 	}
       copy[off] = valid;
     }
+  if (copy)
+    {
+      match = IDENTIFIER_POINTER (get_identifier (copy));
+      free (copy);
+    }
   return match;
 #else
   return ptr;
@@ -919,9 +924,7 @@ validize_symbol_for_target (symtab_node *node)
     {
       symtab->change_decl_assembler_name (decl, get_identifier (name2));
       if (node->lto_file_data)
-	lto_record_renamed_decl (node->lto_file_data, name,
-				 IDENTIFIER_POINTER
-				 (DECL_ASSEMBLER_NAME (decl)));
+	lto_record_renamed_decl (node->lto_file_data, name, name2);
     }
 }
 
@@ -934,12 +937,12 @@ static hash_map<const char *, unsigned> *lto_clone_numbers;
 static bool
 privatize_symbol_name_1 (symtab_node *node, tree decl)
 {
-  const char *name = IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (decl));
+  const char *name0 = IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (decl));
 
-  if (must_not_rename (node, name))
+  if (must_not_rename (node, name0))
     return false;
 
-  name = maybe_rewrite_identifier (name);
+  const char *name = maybe_rewrite_identifier (name0);
   unsigned &clone_number = lto_clone_numbers->get_or_insert (name);
   symtab->change_decl_assembler_name (decl,
 				      clone_function_name (
@@ -947,7 +950,7 @@ privatize_symbol_name_1 (symtab_node *node, tree decl)
   clone_number++;
 
   if (node->lto_file_data)
-    lto_record_renamed_decl (node->lto_file_data, name,
+    lto_record_renamed_decl (node->lto_file_data, name0,
 			     IDENTIFIER_POINTER
 			     (DECL_ASSEMBLER_NAME (decl)));
 
diff --git a/gcc/lto/lto.c b/gcc/lto/lto.c
index 5903f75ac23..cdf75c83613 100644
--- a/gcc/lto/lto.c
+++ b/gcc/lto/lto.c
@@ -424,6 +424,32 @@ lto_wpa_write_files (void)
   timevar_pop (TV_WHOPR_WPA_IO);
 }
 
+/* Create artificial pointers for "omp declare target link" vars.  */
+
+static void
+offload_handle_link_vars (void)
+{
+#ifdef ACCEL_COMPILER
+  varpool_node *var;
+  FOR_EACH_VARIABLE (var)
+    if (lookup_attribute ("omp declare target link",
+			  DECL_ATTRIBUTES (var->decl)))
+      {
+	tree type = build_pointer_type (TREE_TYPE (var->decl));
+	tree link_ptr_var = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+					clone_function_name (var->decl,
+							     "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));
+	DECL_HAS_VALUE_EXPR_P (var->decl) = 1;
+      }
+#endif
+}
+
 /* Perform whole program analysis (WPA) on the callgraph and write out the
    optimization plan.  */
 
@@ -516,6 +542,7 @@ do_whole_program_analysis (void)
      to globals with hidden visibility because they are accessed from multiple
      partitions.  */
   lto_promote_cross_file_statics ();
+  offload_handle_link_vars ();
   if (dump_file)
      dump_end (partition_dump_id, dump_file);
   dump_file = NULL;
@@ -549,32 +576,6 @@ do_whole_program_analysis (void)
     dump_memory_report ("Final");
 }
 
-/* Create artificial pointers for "omp declare target link" vars.  */
-
-static void
-offload_handle_link_vars (void)
-{
-#ifdef ACCEL_COMPILER
-  varpool_node *var;
-  FOR_EACH_VARIABLE (var)
-    if (lookup_attribute ("omp declare target link",
-			  DECL_ATTRIBUTES (var->decl)))
-      {
-	tree type = build_pointer_type (TREE_TYPE (var->decl));
-	tree link_ptr_var = build_decl (UNKNOWN_LOCATION, VAR_DECL,
-					clone_function_name (var->decl,
-							     "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));
-	DECL_HAS_VALUE_EXPR_P (var->decl) = 1;
-      }
-#endif
-}
-
 unsigned int
 lto_option_lang_mask (void)
 {
@@ -641,7 +642,10 @@ lto_main (void)
 
 	  materialize_cgraph ();
 	  if (!flag_ltrans)
-	    lto_promote_statics_nonwpa ();
+	    {
+	      lto_promote_statics_nonwpa ();
+	      offload_handle_link_vars ();
+	    }
 
 	  /* Annotate the CU DIE and mark the early debug phase as finished.  */
 	  debuginfo_early_start ();
diff --git a/libgomp/testsuite/libgomp.c++/target-same-name-2-a.C b/libgomp/testsuite/libgomp.c++/target-same-name-2-a.C
new file mode 100644
index 00000000000..1cff1c8d0c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-same-name-2-a.C
@@ -0,0 +1,50 @@
+/* { dg-skip-if "" { *-*-* } } */
+/* Used by target-same-name-2.c */
+
+#include <complex>
+
+template<typename T>
+int
+test_map ()
+{
+  std::complex<T> a(2, 1), a_check;
+#pragma omp target map(from : a_check)
+  {
+    a_check = a;
+  }
+  if (a == a_check)
+    return 42;
+  return 0;
+}
+
+template<typename T>
+static int
+test_map_static ()
+{
+  std::complex<T> a(-4, 5), a_check;
+#pragma omp target map(from : a_check)
+  {
+    a_check = a;
+  }
+  if (a == a_check)
+    return 441;
+  return 0;
+}
+
+int
+test_a ()
+{
+  int res = test_map<float>();
+  if (res != 42)
+    __builtin_abort ();
+  return res;
+}
+
+int
+test_a2 ()
+{
+  int res = test_map_static<float>();
+  if (res != 441)
+    __builtin_abort ();
+  return res;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-same-name-2-b.C b/libgomp/testsuite/libgomp.c++/target-same-name-2-b.C
new file mode 100644
index 00000000000..31884ba57ce
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-same-name-2-b.C
@@ -0,0 +1,50 @@
+/* { dg-skip-if "" { *-*-* } } */
+/* Used by target-same-name-2.c */
+
+#include <complex>
+
+template<typename T>
+int
+test_map ()
+{
+  std::complex<T> a(2, 1), a_check;
+#pragma omp target map(from : a_check)
+  {
+    a_check = a;
+  }
+  if (a == a_check)
+    return 42;
+  return 0;
+}
+
+template<typename T>
+static int
+test_map_static ()
+{
+  std::complex<T> a(-4, 5), a_check;
+#pragma omp target map(from : a_check)
+  {
+    a_check = a;
+  }
+  if (a == a_check)
+    return 442;
+  return 0;
+}
+
+int
+test_b()
+{
+  int res = test_map<float>();
+  if (res != 42)
+    __builtin_abort ();
+  return res;
+}
+
+int
+test_b2()
+{
+  int res = test_map_static<float>();
+  if (res != 442)
+    __builtin_abort ();
+  return res;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-same-name-2.C b/libgomp/testsuite/libgomp.c++/target-same-name-2.C
new file mode 100644
index 00000000000..e14d435d1ff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-same-name-2.C
@@ -0,0 +1,24 @@
+/* { dg-additional-sources "target-same-name-2-a.C target-same-name-2-b.C" } */
+/* PR middle-end/104285 */
+
+/* Both files create the same symbol, which caused issues
+   in non-host lto1. */
+
+int test_a ();
+int test_a2 ();
+int test_b ();
+int test_b2 ();
+
+int
+main ()
+{
+  if (test_a () != 42)
+    __builtin_abort ();
+  if (test_a2 () != 441)
+    __builtin_abort ();
+  if (test_b () != 42)
+    __builtin_abort ();
+  if (test_b2 () != 442)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1-a.c b/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1-a.c
new file mode 100644
index 00000000000..509c238cf8d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1-a.c
@@ -0,0 +1,60 @@
+/* { dg-skip-if "" { *-*-* } } */
+/* Used by target-same-name-1.c */
+
+static int local_link = 42;
+#pragma omp declare target link(local_link)
+
+int decl_a_link = 123;
+#pragma omp declare target link(decl_a_link)
+
+#pragma omp declare target
+static int __attribute__ ((noinline,noclone))
+foo ()
+{
+  return 5;
+}
+#pragma omp end declare target
+
+static int __attribute__ ((noinline,noclone))
+bar ()
+{
+  int i;
+  #pragma omp target map(from:i)
+    i = foo ();
+  return i;
+}
+
+int
+one () {
+  return bar ();
+}
+
+int
+one_get_inc2_local_link ()
+{
+  int res, res2;
+#pragma omp target map(from: res, res2)
+  {
+    res = local_link;
+    local_link += 2;
+    res2 = local_link;
+  }
+  if (res + 2 != res2)
+    __builtin_abort ();
+  return res;
+}
+
+int
+one_get_inc3_link_a ()
+{
+  int res, res2;
+#pragma omp target map(from: res, res2)
+  {
+    res = decl_a_link;
+    decl_a_link += 3;
+    res2 = decl_a_link;
+  }
+  if (res + 3 != res2)
+    __builtin_abort ();
+  return res;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1-b.c b/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1-b.c
new file mode 100644
index 00000000000..ce008762797
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1-b.c
@@ -0,0 +1,60 @@
+/* { dg-skip-if "" { *-*-* } } */
+/* Used by target-same-name-1.c */
+
+static int local_link = 55;
+#pragma omp declare target link(local_link)
+
+extern int decl_a_link;
+#pragma omp declare target link(decl_a_link)
+
+#pragma omp declare target
+static int __attribute__ ((noinline,noclone))
+foo ()
+{
+  return 7;
+}
+#pragma omp end declare target
+
+static int __attribute__ ((noinline,noclone))
+bar ()
+{
+  int i;
+  #pragma omp target map(from:i)
+    i = foo ();
+  return i;
+}
+
+int
+two () {
+  return bar ();
+}
+
+int
+two_get_inc4_local_link ()
+{
+  int res, res2;
+#pragma omp target map(from: res, res2)
+  {
+    res = local_link;
+    local_link += 4;
+    res2 = local_link;
+  }
+  if (res + 4 != res2)
+    __builtin_abort ();
+  return res;
+}
+
+int
+two_get_inc5_link_a ()
+{
+  int res, res2;
+#pragma omp target map(from: res, res2)
+  {
+    res = decl_a_link;
+    decl_a_link += 5;
+    res2 = decl_a_link;
+  }
+  if (res + 5 != res2)
+    __builtin_abort ();
+  return res;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1.c
new file mode 100644
index 00000000000..b35d8c96ae2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-same-name-1.c
@@ -0,0 +1,46 @@
+/* { dg-additional-sources "target-same-name-1-a.c target-same-name-1-b.c" } */
+/* PR middle-end/104285 */
+
+/* Both files create the same static symbol, which caused issues
+   in non-host lto1. */
+
+int one ();
+int two ();
+int one_get_inc2_local_link ();
+int two_get_inc4_local_link ();
+int one_get_inc3_link_a ();
+int two_get_inc5_link_a ();
+
+int
+main ()
+{
+  if (one () != 5)
+    __builtin_abort ();
+  if (two () != 7)
+    __builtin_abort ();
+
+  if (one_get_inc2_local_link () != 42)
+    __builtin_abort ();
+  if (two_get_inc4_local_link () != 55)
+    __builtin_abort ();
+  if (one_get_inc2_local_link () != 42+2)
+    __builtin_abort ();
+  if (two_get_inc4_local_link () != 55+4)
+    __builtin_abort ();
+
+  if (one_get_inc3_link_a () != 123)
+    __builtin_abort ();
+  if (two_get_inc5_link_a () != 123+3)
+    __builtin_abort ();
+
+/* FIXME: The last call did not increment the global var. */
+/* PR middle-end/105015  */
+#if 0
+  if (one_get_inc3_link_a () != 123+3+5)
+    __builtin_abort ();
+  if (two_get_inc5_link_a () != 123+3+5+3)
+    __builtin_abort ();
+#endif
+
+  return 0;
+}