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: [...]