From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id 55F843858281; Fri, 20 Jan 2023 20:39:47 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 55F843858281 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1674247187; bh=abo23ffPhHcyS7n99P0RQrVNzR/2kJ5AutFCCFAQEbU=; h=From:To:Subject:Date:From; b=Ry/8SGeJP5ZkL3+sGb+n6n5tjdoj3DpSLrxzNv9rXdn8EUjtswZh9ttWmsMI/5bXk 9hloQZikzJ7oJpSP2jl4AxfqTE+QRQRG0stEV0sVuHK5qDtk262/T3L8Kw0k7xpS5E AcO6aGE4YmdoOAXb1g1rt14hRWJ56DgfV5xrCUH4= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] nvptx: Support global constructors/destructors via 'collect2' X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: ea52f1ca16870e4228f8044588b1bf958d4723b0 X-Git-Newrev: fe07b0003bb2092bc34d4bed504be1868b88782d Message-Id: <20230120203947.55F843858281@sourceware.org> Date: Fri, 20 Jan 2023 20:39:47 +0000 (GMT) List-Id: https://gcc.gnu.org/g:fe07b0003bb2092bc34d4bed504be1868b88782d commit fe07b0003bb2092bc34d4bed504be1868b88782d Author: Thomas Schwinge Date: Sun Nov 13 14:19:30 2022 +0100 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. Diff: --- gcc/ChangeLog.omp | 5 ++ gcc/collect2.cc | 4 ++ gcc/config.gcc | 1 + gcc/config/nvptx/nvptx.h | 35 ++++++++++++- gcc/testsuite/ChangeLog.omp | 6 +++ .../gcc.dg/no_profile_instrument_function-attr-1.c | 2 +- gcc/testsuite/lib/target-supports.exp | 3 +- libgcc/ChangeLog.omp | 9 ++++ libgcc/config.host | 2 +- libgcc/config/nvptx/crt0.c | 6 +++ libgcc/config/nvptx/crtstuff.c | 58 ++++++++++++++++++++++ libgcc/config/nvptx/t-nvptx | 15 +++++- 12 files changed, 139 insertions(+), 7 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 127b450644b..ca00bfb48f9 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,5 +1,10 @@ 2023-01-20 Thomas Schwinge + * collect2.cc (write_c_file_glob): Allow for + 'COLLECT2_MAIN_REFERENCE' override. + * config.gcc : Set 'use_collect2=yes'. + * config/nvptx/nvptx.h: Adjust. + * config/nvptx/nvptx.cc (nvptx_assemble_undefined_decl): Notice '__nvptx_stacks', '__nvptx_uni' declarations. (nvptx_file_end): Don't emit duplicate declarations for those. diff --git a/gcc/collect2.cc b/gcc/collect2.cc index d81c7f28f16..945a9ff86dd 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 e6b9c864b0d..9c9365886cf 100644 --- a/gcc/config.gcc +++ b/gcc/config.gcc @@ -2835,6 +2835,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 b81f9f42cd3..d815081147e 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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index c942c34dc70..f35568d83a9 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,5 +1,11 @@ 2023-01-20 Thomas Schwinge + * 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. + * gcc.target/nvptx/softstack-decl-1.c: Make 'dg-do assemble', adjust. * gcc.target/nvptx/uniform-simt-decl-1.c: Likewise. 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 909f8a68479..5b4101cf596 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 e911b298f4c..efd02a61a6b 100644 --- a/gcc/testsuite/lib/target-supports.exp +++ b/gcc/testsuite/lib/target-supports.exp @@ -850,8 +850,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/ChangeLog.omp b/libgcc/ChangeLog.omp index f41bbbb339a..68a99cbe427 100644 --- a/libgcc/ChangeLog.omp +++ b/libgcc/ChangeLog.omp @@ -1,3 +1,12 @@ +2023-01-20 Thomas Schwinge + + * 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. + 2022-11-07 Kwok Cheung Yeung * config/gcn/simd-math/amdgcnmach.h (VECTOR_RETURN): Store value of diff --git a/libgcc/config.host b/libgcc/config.host index 072dd26a276..29b5a2bd7cf 100644 --- a/libgcc/config.host +++ b/libgcc/config.host @@ -1535,7 +1535,7 @@ m32c-*-elf*|m32c-*-rtems*) ;; 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 abf047327ae..860e2bfacad 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 00000000000..0823fc49901 --- /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 ede0bf0f87d..9a0454c3a4d 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