| Message ID | 1adcf079-7ac8-544d-d896-d6d246e689e1@acm.org |
|---|---|
| State | New |
| Headers | show |
On Wed, 25 May 2016, Nathan Sidwell wrote: > This patch reimplements crt0 as C rather than assembly. That means it;ll be > good with 32 bit addresses. I've also completely moved abort and exit into > newlib (a fork on github), which had it's own copies that failed to set the > exit code. Things happened to work because they were never found, being > hidden by the ones provided in crt0.s. You'll need to update your newlib > sources for exit codes to continue to work. Doesn't this break offloading compilation? When you link offloaded code against new newlib, nothing provides a definition of __exitval_ptr (crt0 is only linked in with -mmainkernel), and PTX does not support ELF semantics for undefined weak symbols where they get a null address. OpenACC testcases that call abort() should get a link error after this change. Also note that &__exitval_ptr is always non-null for that reason. > Alex, this should make it simple to add a check in __main for a single thread > soft stack allocation. I imagine a check on a weakly declared __soft_stack > symbol or something? For the above reason weak linkage is not usable in that fashion. Code on my branch simply sets up soft stack in __main regardless of whether it's used. > Or is that something nvptx-run should check for and initialize? The interface between nvptx-run and crt0 would need to change to allow that. Alexander
On 05/25/16 09:28, Alexander Monakov wrote: > On Wed, 25 May 2016, Nathan Sidwell wrote: > >> This patch reimplements crt0 as C rather than assembly. That means it;ll be >> good with 32 bit addresses. I've also completely moved abort and exit into >> newlib (a fork on github), which had it's own copies that failed to set the >> exit code. Things happened to work because they were never found, being >> hidden by the ones provided in crt0.s. You'll need to update your newlib >> sources for exit codes to continue to work. > > Doesn't this break offloading compilation? It won't -- but you're right about .weak not being undefined-is-zero. >> Or is that something nvptx-run should check for and initialize? > > The interface between nvptx-run and crt0 would need to change to allow that. That is exactly what I'm suggesting. nvptx-run looks for a __soft_stack symbol and goes poking at it before launch. nathan
On Wed, 25 May 2016, Nathan Sidwell wrote: > On 05/25/16 09:28, Alexander Monakov wrote: > > On Wed, 25 May 2016, Nathan Sidwell wrote: > > > > > This patch reimplements crt0 as C rather than assembly. That means > > > it;ll be > > > good with 32 bit addresses. I've also completely moved abort and exit > > > into > > > newlib (a fork on github), which had it's own copies that failed to set > > > the > > > exit code. Things happened to work because they were never found, being > > > hidden by the ones provided in crt0.s. You'll need to update your newlib > > > sources for exit codes to continue to work. > > > > Doesn't this break offloading compilation? > > It won't -- but you're right about .weak not being undefined-is-zero. With today's trunk and newlib, if I run make check-target-libgomp RUNTESTFLAGS=c.exp=abort-2.c in the output I see: Running ../../../../gcc/libgomp/testsuite/libgomp.oacc-c/c.exp ... FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/abort-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 (test for excess errors) FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/abort-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 (test for excess errors) FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/abort-2.c -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -O2 (test for excess errors) and libgomp.log has: unresolved symbol __exitval_ptr collect2: error: ld returned 1 exit status mkoffload: fatal error: /home/am/ww/i/bin//x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status compilation terminated. lto-wrapper: fatal error: /home/am/ww/i/libexec/gcc/x86_64-pc-linux-gnu/7.0.0//accel/nvptx-none/mkoffload returned 1 exit status compilation terminated. /usr/bin/ld: lto-wrapper failed collect2: error: ld returned 1 exit status compiler exited with status 1 Am I missing anything? > > > Or is that something nvptx-run should check for and initialize? > > > > The interface between nvptx-run and crt0 would need to change to allow that. > > That is exactly what I'm suggesting. nvptx-run looks for a __soft_stack > symbol and goes poking at it before launch. It is possible, but it seems it's enough to set up soft stacks under #ifdef __nvptx_softstack__ in __main. Is that fine? Alexander
On 05/25/16 11:49, Alexander Monakov wrote: > On Wed, 25 May 2016, Nathan Sidwell wrote: > With today's trunk and newlib, if I run > > > unresolved symbol __exitval_ptr Is should work now, just pushed a patch to newlib. PTX appears to accept '.extern .weak ...', but that has the same semantics as '.extern ...', which IMHO is a bug. '.extern .weak' doesn't mean anything special. Working on a GCC patch to stop us emitting it. > It is possible, but it seems it's enough to set up soft stacks under #ifdef > __nvptx_softstack__ in __main. Is that fine? That's fine -- I'd forgotten there was a #define to check. The whole point of reimplementing crt0 in C was to make that kind of thing easier! nathan
On 05/26/16 10:36, Nathan Sidwell wrote: > Ib. PTX appears to accept > '.extern .weak ...', but that has the same semantics as '.extern ...', which > IMHO is a bug. '.extern .weak' doesn't mean anything special. Working on a > GCC patch to stop us emitting it. Pah, I'd misremembered what we emitted. Still working on a patch to stop weakrefs ... nathan
2016-05-25 Nathan Sidwell <nathan@acm.org> libgcc/ * config/nvptx/crt0.s: Delete. * config/nvptx/crt0.c: New. * t-nvptx: Update. gcc/testsuite/ * gcc.c-torture/execute/921110-1.c: Fix abort decl. Index: gcc/testsuite/gcc.c-torture/execute/921110-1.c =================================================================== --- gcc/testsuite/gcc.c-torture/execute/921110-1.c (revision 236531) +++ gcc/testsuite/gcc.c-torture/execute/921110-1.c (working copy) @@ -1,7 +1,8 @@ -extern int abort(); -typedef int (*frob)(); +extern void abort(void); +typedef void (*frob)(); frob f[] = {abort}; -main() + +int main(void) { exit(0); } Index: libgcc/config/nvptx/crt0.c =================================================================== --- libgcc/config/nvptx/crt0.c (nonexistent) +++ libgcc/config/nvptx/crt0.c (working copy) @@ -0,0 +1,37 @@ +/* Copyright (C) 2014-2016 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 + <http://www.gnu.org/licenses/>. */ + +int *__exitval_ptr; + +extern void __attribute__((noreturn)) exit (int status); +extern int main (int, void **); + +void __attribute__((kernel)) +__main (int *rval_ptr, int argc, void **argv) +{ + __exitval_ptr = rval_ptr; + /* Store something non-zero, so the host knows something went wrong, + if we fail to reach exit properly. */ + if (rval_ptr) + *rval_ptr = 255; + + exit (main (argc, argv)); +} Index: libgcc/config/nvptx/crt0.s =================================================================== --- libgcc/config/nvptx/crt0.s (revision 236531) +++ libgcc/config/nvptx/crt0.s (nonexistent) @@ -1,45 +0,0 @@ - .version 3.1 - .target sm_30 - .address_size 64 - -.global .u64 %__exitval; -// BEGIN GLOBAL FUNCTION DEF: abort -.visible .func abort -{ - .reg .u64 %rd1; - ld.global.u64 %rd1,[%__exitval]; - st.u32 [%rd1], 255; - exit; -} -// BEGIN GLOBAL FUNCTION DEF: exit -.visible .func exit (.param .u32 %arg) -{ - .reg .u64 %rd1; - .reg .u32 %val; - ld.param.u32 %val,[%arg]; - ld.global.u64 %rd1,[%__exitval]; - st.u32 [%rd1], %val; - exit; -} - -.extern .func (.param.u32 retval) main (.param.u32 argc, .param.u64 argv); - -.visible .entry __main (.param .u64 __retval, .param.u32 __argc, .param.u64 __argv) -{ - .reg .u32 %r<3>; - .reg .u64 %rd<3>; - .param.u32 %argc; - .param.u64 %argp; - .param.u32 %mainret; - ld.param.u64 %rd0, [__retval]; - st.global.u64 [%__exitval], %rd0; - - ld.param.u32 %r1, [__argc]; - ld.param.u64 %rd1, [__argv]; - st.param.u32 [%argc], %r1; - st.param.u64 [%argp], %rd1; - call.uni (%mainret), main, (%argc, %argp); - ld.param.u32 %r1,[%mainret]; - st.s32 [%rd0], %r1; - exit; -} Index: libgcc/config/nvptx/t-nvptx =================================================================== --- libgcc/config/nvptx/t-nvptx (revision 236531) +++ libgcc/config/nvptx/t-nvptx (working copy) @@ -6,8 +6,8 @@ LIB2ADD=$(srcdir)/config/nvptx/malloc.as LIB2ADDEH= LIB2FUNCS_EXCLUDE=__main -crt0.o: $(srcdir)/config/nvptx/crt0.s - cp $< $@ +crt0.o: $(srcdir)/config/nvptx/crt0.c + $(crt_compile) -c $< # Prevent building "advanced" stuff (for example, gcov support). We don't # support it, and it may cause the build to fail, because of alloca usage, for