Hi Jakub, hi Richi,
On 22.03.22 15:27, Jakub Jelinek wrote:
On Tue, Mar 22, 2022 at 01:56:27PM +0100, Tobias Burnus wrote:
name = lto_get_decl_name_mapping (file_data, name);
[...]
+ name = lto_get_decl_name_mapping (file_data, name);
This looks weird. IMHO we should make sure that the hash table contains
always the final names, so that we don't need to call it twice or even more
times.
I concur – and I have changed it (see attached patch). And it still passes the
testsuite.
(It came about because I started reverse – and in
'cgraph_node::get_untransformed_body',
it happened that the second map was already present + I added the second lookup.
Extending the tests, I found that I also had to change privatize_symbol_name_1.
Thus,
I added the additional mapping there. However, it seems to be enough to only
change
privatize_symbol_name_1 to track the double renaming in one step.
That's what the updated patch does (+ unchanged changes for linkptr + the
hash-key part).
The lto.cc changes LGTM. The rest I'm unfortunately not too familiar with,
Richi or Honza, could you please have a look?
(^ to be done by Honza or Richi)
This is the same thing again as above, right?
Yes. (Now also removed.)
Thanks for the first review and for nagging about the rename tracking.
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
LTO: Fixes for renaming issues with offload/OpenMP [PR104285]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.
gcc/lto/lto-partition.cc | 17 +++---
gcc/lto/lto.cc | 58 +++++++++++----------
.../testsuite/libgomp.c++/target-same-name-2-a.C | 50 ++++++++++++++++++
.../testsuite/libgomp.c++/target-same-name-2-b.C | 50 ++++++++++++++++++
libgomp/testsuite/libgomp.c++/target-same-name-2.C | 24 +++++++++
.../libgomp.c-c++-common/target-same-name-1-a.c | 60 ++++++++++++++++++++++
.../libgomp.c-c++-common/target-same-name-1-b.c | 60 ++++++++++++++++++++++
.../libgomp.c-c++-common/target-same-name-1.c | 46 +++++++++++++++++
8 files changed, 331 insertions(+), 34 deletions(-)
diff --git a/gcc/lto/lto-partition.cc b/gcc/lto/lto-partition.cc
index 062fd033ecb..ebb9c3abe12 100644
--- 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 98c336a152b..31b0c1862f7 100644
--- 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
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;
+}
