LTO: Fixes for renaming issues with offload/OpenMP [PR104285]
Commit Message
This patch fixes some same-local-name issues with offloading.
The first issue can also occur with -flto with -foffload=disable (I think).
Otherwise, all three issues apply to the non-host lto1 run
(defined ACCEL_COMPILER).
First, for omp declare link(var), a variable is declared
based on DECL_NAME (var) by appending "$linkptr"
(→ offload_handle_link_vars) – and saved as VALUE_EXPR.
The problem is that this call happens before the static
variables are renamed in lto_promote_statics_nonwpa
or lto_promote_cross_file_statics.
Solution: Call offload_handle_link_vars right after those
static-name handlings. (Plus move offload_handle_link_vars
before the first caller.)
* * *
Secondly, maybe_rewrite_identifier: If the ACCEL_COMPILER
is more restricted than the host compiler, . (dot) or '$'
have to be replaced. Thus, this function is only active for
#if defined ACCEL_COMPILER && (defined NO_DOT_IN_LABEL || defined NO_DOLLAR_IN_LABEL)
This function created a new string (xstrdup) and returned it as it.
That works kind of okay, but not for duplicated names:
unsigned &clone_number = lto_clone_numbers->get_or_insert (name);
The underlying hash function uses the pointer address and not the
string value to return the clone_number – resulting in '0' for
xstrdup strings.
Solution: maybe_rewrite_identifier now uses IDENTIFIER_POINTER (get_identifier (...))
to get a unique pointer.
* * *
Last issue: When a function name was changed because of $. issues (e.g. 'func._omp_fn.0'
to 'func$_omp_fn$0') and then the function name was changed because of static-name
collisions< ('func$_omp_fn$0' → 'func$_omp_fn$0$ltrans_0').
In this case, we need to go back to the original name to obtain the function body.
The second renaming has been properly tracked already - but the first one hasn't.
This patch now also tracks the first renaming + calls lto_get_decl_name_mapping a
second time. Instead of using
+#if defined ACCEL_COMPILER && (defined NO_DOT_IN_LABEL || defined NO_DOLLAR_IN_LABEL)
+ name = lto_get_decl_name_mapping (file_data, name);
+#endif
the code could have also used the call unconditionally as it returns the passed
argument when no mapping could be found. The condition mirrors the one for
maybe_rewrite_identifier.
The idea was to have an optimization + annotate the code a bit. An alternative would
be to remove the '#' lines, possibly replacing the first line
by '/* Handle maybe_rewrite_identifier renames. */'.
OK for current trunk (GCC12)? Alternatively, OK for GCC 13 Stage 1?
Tobias
PS: After having finished this patch for the PR, I did encounter an issue with
the omptests testsuite and linkptr (see 'first' in this email); with this patch
the omptests test case passes. But while writing a testcase for linkptr, I did
run into https://gcc.gnu.org/PR105015 (see PR or #if 0 code in one of the testcase
files in the attached patch).
-----------------
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
Comments
On Tue, Mar 22, 2022 at 01:56:27PM +0100, Tobias Burnus wrote:
> --- a/gcc/cgraph.cc
> +++ b/gcc/cgraph.cc
> @@ -3980,6 +3980,9 @@ cgraph_node::get_untransformed_body ()
>
> /* We may have renamed the declaration, e.g., a static function. */
> name = lto_get_decl_name_mapping (file_data, name);
> +#if defined ACCEL_COMPILER && (defined NO_DOT_IN_LABEL || defined NO_DOLLAR_IN_LABEL)
> + name = lto_get_decl_name_mapping (file_data, name);
> +#endif
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.
> --- 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 ();
The lto.cc changes LGTM. The rest I'm unfortunately not too familiar with,
Richi or Honza, could you please have a look?
> --- a/gcc/varpool.cc
> +++ b/gcc/varpool.cc
> @@ -294,6 +294,9 @@ varpool_node::get_constructor (void)
>
> /* We may have renamed the declaration, e.g., a static function. */
> name = lto_get_decl_name_mapping (file_data, name);
> +#if defined ACCEL_COMPILER && (defined NO_DOT_IN_LABEL || defined NO_DOLLAR_IN_LABEL)
> + name = lto_get_decl_name_mapping (file_data, name);
> +#endif
> struct lto_in_decl_state *decl_state
> = lto_get_function_in_decl_state (file_data, decl);
This is the same thing again as above, right?
Jakub
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
On Tue, 22 Mar 2022, Tobias Burnus wrote:
> 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.
OK.
Thanks,
Richard.
LTO: Fixes for renaming issues with offload/OpenMP [PR104285]LTO: Fixes for renaming issues with offload/OpenMP [PR104285]
gcc/ChangeLog:
PR middle-end/104285
* cgraph.cc (cgraph_node::get_untransformed_body): Call
lto_get_decl_name_mapping again to handle NO_DOT/NO_DOLLAR renames
for ACCEL_COMPILER.
* varpool.cc (varpool_node::get_constructor): Likewise.
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/cgraph.cc | 3 ++
gcc/lto/lto-partition.cc | 18 ++++---
gcc/lto/lto.cc | 58 +++++++++++----------
gcc/varpool.cc | 3 ++
.../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 +++++++++++++++++
10 files changed, 339 insertions(+), 33 deletions(-)
@@ -3980,6 +3980,9 @@ cgraph_node::get_untransformed_body ()
/* We may have renamed the declaration, e.g., a static function. */
name = lto_get_decl_name_mapping (file_data, name);
+#if defined ACCEL_COMPILER && (defined NO_DOT_IN_LABEL || defined NO_DOLLAR_IN_LABEL)
+ name = lto_get_decl_name_mapping (file_data, name);
+#endif
struct lto_in_decl_state *decl_state
= lto_get_function_in_decl_state (file_data, decl);
@@ -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,18 +939,21 @@ 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 (
name, "lto_priv", clone_number));
clone_number++;
+ if (node->lto_file_data && name0 != name)
+ lto_record_renamed_decl (node->lto_file_data, name0, name);
+
if (node->lto_file_data)
lto_record_renamed_decl (node->lto_file_data, name,
IDENTIFIER_POINTER
@@ -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 ();
@@ -294,6 +294,9 @@ varpool_node::get_constructor (void)
/* We may have renamed the declaration, e.g., a static function. */
name = lto_get_decl_name_mapping (file_data, name);
+#if defined ACCEL_COMPILER && (defined NO_DOT_IN_LABEL || defined NO_DOLLAR_IN_LABEL)
+ name = lto_get_decl_name_mapping (file_data, name);
+#endif
struct lto_in_decl_state *decl_state
= lto_get_function_in_decl_state (file_data, decl);
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}