From patchwork Wed Jan 11 11:48:31 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 62942 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 C416A3858289 for ; Wed, 11 Jan 2023 11:49:04 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 076B13858C52 for ; Wed, 11 Jan 2023 11:48:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 076B13858C52 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.96,315,1665475200"; d="scan'208,223";a="93064443" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 11 Jan 2023 03:48:46 -0800 IronPort-SDR: JwNIE0lMCbIXOzkmFXdUBmzlb42R3CAS4hktnS4KK/9BcLTN4oIaHRlsWzAvg9wxAbbozhdk74 1BAR5+QmbE2j42Y2vLPupxIsyyGFEJn6ZB0f5zwLQ59Scd3Ez3bqQXxwUd3COJcDKyJ8kaVHLB ixUJ19YE14mUEsX9AjtcU2rUxkdMcrIW9ryL/vWY9G661qTinQLBhiU9tcRMRI6OL7cH1/jRs8 nDMLUsZkocCCn3ZdKZOddypRUxqsBX8EvKQwrPlOrH08s6sPNnMu+EErrGRe9qBmDhxgcCfAUt jhU= From: Thomas Schwinge To: , Tom de Vries Subject: [PING^2] nvptx: Support global constructors/destructors via 'collect2' In-Reply-To: <87fsdacxi0.fsf@euler.schwinge.homeip.net> References: <878rjqaku5.fsf@dem-tschwing-1.ger.mentorg.com> <87y1rq7wt4.fsf@dem-tschwing-1.ger.mentorg.com> <87fsdacxi0.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Wed, 11 Jan 2023 12:48:31 +0100 Message-ID: <87h6wxl2b4.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-11.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) 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" Hi! Ping. Grüße Thomas On 2022-12-20T09:03:51+0100, I wrote: > Hi! > > Ping. > > > Minor change in the attached > "nvptx: Support global constructors/destructors via 'collect2'": for > 'atexit', add '#include ' to 'libgcc/config/nvptx/crt0.c'. > > > Grüße > Thomas > > > On 2022-12-02T14:35:35+0100, I wrote: >> Hi! >> >> On 2022-12-01T22:13:38+0100, I wrote: >>> I'm working on support for global constructors/destructors with >>> GCC/nvptx >> >> See "nvptx: Support global constructors/destructors via 'collect2'" >> attached; OK to push? (... with 'gcc/doc/install.texi' accordingly >> updated once >> "'nm'" and newlib >> >> "nvptx: Implement '_exit' instead of 'exit'" have been merged; any >> comments to those?) >> >> Per my quick scanning of 'gcc/config.gcc' history, for more than two >> decades, there was a clear trend to remove 'use_collect2=yes' >> configurations; now finally a new one is being added -- making sure we're >> not slowly dispensing with the need for the early 1990s piece of work >> that 'gcc/collect2*' is... ;'-P >> >> >> Grüße >> Thomas ----------------- 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 From 0e7cf5a9f83c3a82eafa126886e5d92651bfbb30 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Sun, 13 Nov 2022 14:19:30 +0100 Subject: [PATCH] nvptx: Support global constructors/destructors via 'collect2' The function attributes 'constructor', 'destructor', and 'init_priority' now work, as do the C++ features making use of this. Test cases with effective target 'global_constructor' and 'init_priority' now generally work, and 'check-gcc-c++' test results greatly improve; no more "sorry, unimplemented: global constructors not supported on this target". This depends on "'nm'" generally, and for global destructors support: newlib "nvptx: Implement '_exit' instead of 'exit'". gcc/ * collect2.cc (write_c_file_glob): Allow for 'COLLECT2_MAIN_REFERENCE' override. * config.gcc : Set 'use_collect2=yes'. * config/nvptx/nvptx.h: Adjust. gcc/testsuite/ * gcc.dg/no_profile_instrument_function-attr-1.c: GCC/nvptx is 'NO_DOT_IN_LABEL' but not 'NO_DOLLAR_IN_LABEL', so '$' may apper in identifiers. * lib/target-supports.exp (check_effective_target_global_constructor): Enable for nvptx. libgcc/ * config.host : Add 'crtbegin.o', 'crtend.o' to 'extra_parts'. * config/nvptx/crt0.c: Invoke '__do_global_ctors', '__do_global_dtors'. * config/nvptx/crtstuff.c: New. * config/nvptx/t-nvptx: Adjust. --- gcc/collect2.cc | 4 ++ gcc/config.gcc | 1 + gcc/config/nvptx/nvptx.h | 35 ++++++++++- .../no_profile_instrument_function-attr-1.c | 2 +- gcc/testsuite/lib/target-supports.exp | 3 +- libgcc/config.host | 2 +- libgcc/config/nvptx/crt0.c | 6 ++ libgcc/config/nvptx/crtstuff.c | 58 +++++++++++++++++++ libgcc/config/nvptx/t-nvptx | 15 ++++- 9 files changed, 119 insertions(+), 7 deletions(-) create mode 100644 libgcc/config/nvptx/crtstuff.c diff --git a/gcc/collect2.cc b/gcc/collect2.cc index d81c7f28f16a..945a9ff86dda 100644 --- a/gcc/collect2.cc +++ b/gcc/collect2.cc @@ -2238,8 +2238,12 @@ write_c_file_glob (FILE *stream, const char *name ATTRIBUTE_UNUSED) fprintf (stream, "\tdereg_frame,\n"); fprintf (stream, "\t0\n};\n\n"); +# ifdef COLLECT2_MAIN_REFERENCE + fprintf (stream, "%s\n\n", COLLECT2_MAIN_REFERENCE); +# else fprintf (stream, "extern entry_pt %s;\n", NAME__MAIN); fprintf (stream, "entry_pt *__main_reference = %s;\n\n", NAME__MAIN); +# endif } #endif /* ! LD_INIT_SWITCH */ diff --git a/gcc/config.gcc b/gcc/config.gcc index 951902338205..fec67d7b6e40 100644 --- a/gcc/config.gcc +++ b/gcc/config.gcc @@ -2784,6 +2784,7 @@ nvptx-*) tm_file="${tm_file} newlib-stdint.h" use_gcc_stdint=wrap tmake_file="nvptx/t-nvptx" + use_collect2=yes if test x$enable_as_accelerator = xyes; then extra_programs="${extra_programs} mkoffload\$(exeext)" tm_file="${tm_file} nvptx/offload.h" diff --git a/gcc/config/nvptx/nvptx.h b/gcc/config/nvptx/nvptx.h index dc676dcb5fc5..235c1e4d99d5 100644 --- a/gcc/config/nvptx/nvptx.h +++ b/gcc/config/nvptx/nvptx.h @@ -35,7 +35,39 @@ '../../gcc.cc:asm_options', 'HAVE_GNU_AS'. */ #define ASM_SPEC "%{v}" -#define STARTFILE_SPEC "%{mmainkernel:crt0.o%s}" +#define STARTFILE_SPEC \ + STARTFILE_SPEC_MMAINKERNEL \ + " " STARTFILE_SPEC_CDTOR + +#define ENDFILE_SPEC \ + ENDFILE_SPEC_CDTOR + +#define STARTFILE_SPEC_MMAINKERNEL "%{mmainkernel:crt0.o%s}" + +/* Support for global constructors/destructors is implemented via + 'collect2' and the following helpers. */ + +#define STARTFILE_SPEC_CDTOR "crtbegin.o%s" + +#define ENDFILE_SPEC_CDTOR "crtend.o%s" + +/* nvptx does its own wrapping of 'main' + (see 'libgcc/config/nvptx/crt0.c:__main'). */ +#define HAS_INIT_SECTION + +/* For example with old Nvidia Tesla K20c, Driver Version: 361.93.02, the + function pointers stored in the '__CTOR_LIST__', '__DTOR_LIST__' arrays + evidently evaluate to NULL in JIT compilation. Avoiding the use of + assembler names ('write_list_with_asm') doesn't help, but defining a dummy + function next to the arrays apparently does work around this issue... + + The default '__main_reference' synthesized by 'collect2' refers to our + 'crt0.o' '__main' function with incompatible signature: + + error : Function '__main' not declared __global__ in all source files + + Address both these issues via 'COLLECT2_MAIN_REFERENCE'. */ +#define COLLECT2_MAIN_REFERENCE "__attribute__((unused)) static void dummy () {}" #define TARGET_CPU_CPP_BUILTINS() nvptx_cpu_cpp_builtins () @@ -348,7 +380,6 @@ struct GTY(()) machine_function #define MOVE_MAX 8 #define MOVE_RATIO(SPEED) 4 #define FUNCTION_MODE QImode -#define HAS_INIT_SECTION 1 /* The C++ front end insists to link against libstdc++ -- which we don't build. Tell it to instead link against the innocuous libgcc. */ diff --git a/gcc/testsuite/gcc.dg/no_profile_instrument_function-attr-1.c b/gcc/testsuite/gcc.dg/no_profile_instrument_function-attr-1.c index 909f8a684791..5b4101cf596d 100644 --- a/gcc/testsuite/gcc.dg/no_profile_instrument_function-attr-1.c +++ b/gcc/testsuite/gcc.dg/no_profile_instrument_function-attr-1.c @@ -18,7 +18,7 @@ int main () return foo (); } -/* { dg-final { scan-tree-dump-times "__gcov0\[._\]main.* = PROF_edge_counter" 1 "optimized"} } */ +/* { dg-final { scan-tree-dump-times "__gcov0\[$._\]main.* = PROF_edge_counter" 1 "optimized"} } */ /* { dg-final { scan-tree-dump-times "__gcov_indirect_call_profiler_v" 1 "optimized" } } */ /* { dg-final { scan-tree-dump-times "__gcov_time_profiler_counter = " 1 "optimized" } } */ /* { dg-final { scan-tree-dump-times "__gcov_init" 1 "optimized" } } */ diff --git a/gcc/testsuite/lib/target-supports.exp b/gcc/testsuite/lib/target-supports.exp index ea06e21c3a14..b1b1c5b36bc2 100644 --- a/gcc/testsuite/lib/target-supports.exp +++ b/gcc/testsuite/lib/target-supports.exp @@ -907,8 +907,7 @@ proc check_effective_target_nonlocal_goto {} { # Return 1 if global constructors are supported, 0 otherwise. proc check_effective_target_global_constructor {} { - if { [istarget nvptx-*-*] - || [istarget bpf-*-*] } { + if { [istarget bpf-*-*] } { return 0 } return 1 diff --git a/libgcc/config.host b/libgcc/config.host index eb23abe89f5e..25072f41860c 100644 --- a/libgcc/config.host +++ b/libgcc/config.host @@ -1499,7 +1499,7 @@ m32c-*-elf*) ;; nvptx-*) tmake_file="$tmake_file nvptx/t-nvptx" - extra_parts="crt0.o" + extra_parts="crt0.o crtbegin.o crtend.o" ;; *) echo "*** Configuration ${host} not supported" 1>&2 diff --git a/libgcc/config/nvptx/crt0.c b/libgcc/config/nvptx/crt0.c index abf047327ae7..860e2bfacadd 100644 --- a/libgcc/config/nvptx/crt0.c +++ b/libgcc/config/nvptx/crt0.c @@ -19,6 +19,9 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see . */ +#include +#include "gbl-ctors.h" + int *__exitval_ptr; extern void __attribute__((noreturn)) exit (int status); @@ -47,5 +50,8 @@ __main (int *rval_ptr, int argc, void **argv) __nvptx_stacks[0] = stack + sizeof stack; __nvptx_uni[0] = 0; + __do_global_ctors (); + atexit (__do_global_dtors); + exit (main (argc, argv)); } diff --git a/libgcc/config/nvptx/crtstuff.c b/libgcc/config/nvptx/crtstuff.c new file mode 100644 index 000000000000..0823fc499019 --- /dev/null +++ b/libgcc/config/nvptx/crtstuff.c @@ -0,0 +1,58 @@ +/* Copyright (C) 2022 Free Software Foundation, Inc. + + This file is free software; you can redistribute it and/or modify it + under the terms of the GNU General Public License as published by the + Free Software Foundation; either version 3, or (at your option) any + later version. + + This file is distributed in the hope that it will be useful, but + WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + General Public License for more details. + + Under Section 7 of GPL version 3, you are granted additional + permissions described in the GCC Runtime Library Exception, version + 3.1, as published by the Free Software Foundation. + + You should have received a copy of the GNU General Public License and + a copy of the GCC Runtime Library Exception along with this program; + see the files COPYING3 and COPYING.RUNTIME respectively. If not, see + . */ + +#include "gbl-ctors.h" + +/* The common 'crtstuff.c' doesn't quite provide what we need, so we roll our + own. + + There's no technical reason in this configuration here to split the two + functions '__do_global_ctors' and '__do_global_ctors' into two separate + files (via 'CRT_BEGIN' and 'CRT_END'): 'crtbegin.o' and 'crtend.o', but we + do so anyway, for symmetry with other configurations. */ + +#ifdef CRT_BEGIN + +void +__do_global_ctors (void) +{ + DO_GLOBAL_CTORS_BODY; +} + +#elif defined(CRT_END) /* ! CRT_BEGIN */ + +void +__do_global_dtors (void) +{ + /* In this configuration here, there's no way that "this routine is run more + than once [...] when exit is called recursively": for nvptx target, the + call to '__do_global_dtors' is registered via 'atexit', which doesn't + re-enter a function already run. + Therefore, we do *not* "arrange to remember where in the list we left off + processing". */ + func_ptr *p; + for (p = __DTOR_LIST__ + 1; *p; ) + (*p++) (); +} + +#else /* ! CRT_BEGIN && ! CRT_END */ +#error "One of CRT_BEGIN or CRT_END must be defined." +#endif diff --git a/libgcc/config/nvptx/t-nvptx b/libgcc/config/nvptx/t-nvptx index ede0bf0f87dd..9a0454c3a4d0 100644 --- a/libgcc/config/nvptx/t-nvptx +++ b/libgcc/config/nvptx/t-nvptx @@ -3,7 +3,7 @@ LIB2ADD=$(srcdir)/config/nvptx/reduction.c \ $(srcdir)/config/nvptx/atomic.c LIB2ADDEH= -LIB2FUNCS_EXCLUDE=__main +LIB2FUNCS_EXCLUDE= crt0.o: $(srcdir)/config/nvptx/crt0.c $(crt_compile) -c $< @@ -12,3 +12,16 @@ crt0.o: $(srcdir)/config/nvptx/crt0.c # support it, and it may cause the build to fail, because of alloca usage, for # example. INHIBIT_LIBC_CFLAGS = -Dinhibit_libc + +# Support for global constructors/destructors is implemented via +# 'collect2' and the following helpers. + +LIB2FUNCS_EXCLUDE += __main + +CUSTOM_CRTSTUFF = yes + +crtbegin.o: $(srcdir)/config/nvptx/crtstuff.c + $(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $< -DCRT_BEGIN + +crtend.o: $(srcdir)/config/nvptx/crtstuff.c + $(crt_compile) $(CRTSTUFF_T_CFLAGS) -c $< -DCRT_END -- 2.35.1