public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] LTO: Fixes for renaming issues with offload/OpenMP [PR104285]
@ 2022-03-23 10:20 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-03-23 10:20 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:d41e8d071577fafe18e7e8046978a2ef9a0fa711
commit d41e8d071577fafe18e7e8046978a2ef9a0fa711
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Wed Mar 23 09:44:39 2022 +0100
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)
Diff:
---
gcc/lto/lto-partition.c | 17 +++---
gcc/lto/lto.c | 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.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;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-03-23 10:20 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-03-23 10:20 [gcc/devel/omp/gcc-11] LTO: Fixes for renaming issues with offload/OpenMP [PR104285] Tobias Burnus
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).