https://gcc.gnu.org/bugzilla/show_bug.cgi?id=112858
--- Comment #1 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Indeed in 'build-gcc/nvptx-none/libstdc++-v3/libsupc++/atexit_thread.o' I see:
// BEGIN GLOBAL FUNCTION DECL: __cxa_thread_atexit_impl
.extern .func (.param .u32 %value_out) __cxa_thread_atexit_impl (.param
.u64 %in_ar0, .param .u64 %in_ar1, .param .u64 %in_ar2);
That is, '.extern' instead of '.weak' linking directive, huh.
..., but still doing the NULL check:
[...]
.reg .u64 %r29;
.reg .pred %r30;
[...]
mov.u64 %r29,__cxa_thread_atexit_impl;
setp.eq.u64 %r30,%r29,0;
@ %r30 bra $L9;
.loc 2 156 37
{
.param .u32 %value_in;
.param .u64 %out_arg1;
st.param.u64 [%out_arg1],%r26;
.param .u64 %out_arg2;
st.param.u64 [%out_arg2],%r27;
.param .u64 %out_arg3;
st.param.u64 [%out_arg3],%r28;
call (%value_in),__cxa_thread_atexit_impl,(%out_arg1,%out_arg2,%out_arg3);
ld.param.u32 %r34,[%value_in];
}
mov.u32 %r25,%r34;
.loc 2 156 59
bra $L8;
$L9:
[...]