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

gcc/lto/ChangeLog:

	PR middle-end/104285
	* lto-partition.cc (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.cc (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.

3 files changed, 166 insertions, 0 deletions

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 0000000..509c238
--- /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 0000000..ce00876
--- /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 0000000..b35d8c9
--- /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;
+}