From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id D8EC63858405; Wed, 23 Mar 2022 10:20:02 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org D8EC63858405 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] LTO: Fixes for renaming issues with offload/OpenMP [PR104285] X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: 1073e7691804f43993cd9f8446f8ebda9004ad98 X-Git-Newrev: d41e8d071577fafe18e7e8046978a2ef9a0fa711 Message-Id: <20220323102002.D8EC63858405@sourceware.org> Date: Wed, 23 Mar 2022 10:20:02 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Wed, 23 Mar 2022 10:20:03 -0000 https://gcc.gnu.org/g:d41e8d071577fafe18e7e8046978a2ef9a0fa711 commit d41e8d071577fafe18e7e8046978a2ef9a0fa711 Author: Tobias Burnus 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 *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 + +template +int +test_map () +{ + std::complex a(2, 1), a_check; +#pragma omp target map(from : a_check) + { + a_check = a; + } + if (a == a_check) + return 42; + return 0; +} + +template +static int +test_map_static () +{ + std::complex 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(); + if (res != 42) + __builtin_abort (); + return res; +} + +int +test_a2 () +{ + int res = test_map_static(); + 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 + +template +int +test_map () +{ + std::complex a(2, 1), a_check; +#pragma omp target map(from : a_check) + { + a_check = a; + } + if (a == a_check) + return 42; + return 0; +} + +template +static int +test_map_static () +{ + std::complex 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(); + if (res != 42) + __builtin_abort (); + return res; +} + +int +test_b2() +{ + int res = test_map_static(); + 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; +}