From patchwork Tue Mar 22 12:56:27 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 52215 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id AC2293899427 for ; Tue, 22 Mar 2022 12:56:54 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id AB57D3857C4A for ; Tue, 22 Mar 2022 12:56:35 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org AB57D3857C4A Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.90,201,1643702400"; d="diff'?scan'208";a="76035501" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 22 Mar 2022 04:56:34 -0800 IronPort-SDR: SK6OFZdr5KH+Zl1fFMXEoUELkO+FJ/6Q38bRBx8SU4OJ1ILOuNiAHxaGs69G2TcIcBtvBNtP2L Hfd1N9wJIkqkimLH+Re3LhsD2qIdU9JdujiEDIGw8db7kCfvcSVhpzl1BtloRXyhOjmytLPzMv w5cD/kMSgx1Odq2SwtFx41KB/PEujCgXGzNe2QGdVQW9Z4S+sNDymMghHjKmaW6IBVaFT3oWRO qusEUdQdGCMYjghcOeglAEoRAhQHhCqQwnawcP+wuWMf6p7//Jt6larVepiuGedMzo++95SpPm Q14= Message-ID: Date: Tue, 22 Mar 2022 13:56:27 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:91.0) Gecko/20100101 Thunderbird/91.7.0 Content-Language: en-US To: gcc-patches , Jakub Jelinek , Richard Biener From: Tobias Burnus Subject: [Patch] LTO: Fixes for renaming issues with offload/OpenMP [PR104285] X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-11.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" 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 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(-) diff --git a/gcc/cgraph.cc b/gcc/cgraph.cc index b923a59ab0c..8b6215662c9 100644 --- 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 struct lto_in_decl_state *decl_state = lto_get_function_in_decl_state (file_data, decl); diff --git a/gcc/lto/lto-partition.cc b/gcc/lto/lto-partition.cc index 062fd033ecb..51323117195 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,18 +939,21 @@ 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 ( 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 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/gcc/varpool.cc b/gcc/varpool.cc index bfd17f1250c..70a1a2ddeee 100644 --- 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); 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; +}