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