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.

diff --git a/gcc/lto/lto-partition.cc b/gcc/lto/lto-partition.cc
index 062fd03..ebb9c3a 100644 (file)
--- a/gcc/lto/lto-partition.cc
+++ b/gcc/lto/lto-partition.cc
@@ -898,6 +898,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;
@@ -921,9 +926,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);
     }
 }
 
@@ -936,12 +939,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 (
@@ -949,7 +952,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.cc b/gcc/lto/lto.cc
index 98c336a..31b0c18 100644 (file)
--- a/gcc/lto/lto.cc
+++ b/gcc/lto/lto.cc
@@ -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 (file)
index 0000000..1cff1c8
--- /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 (file)
index 0000000..31884ba
--- /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 (file)
index 0000000..e14d435
--- /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 (file)
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 (file)
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 (file)
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;
+}