From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-bugzilla@gcc.gnu.org>
Received: by sourceware.org (Postfix, from userid 48)
	id 9E65038618B3; Tue,  5 Dec 2023 08:07:09 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 9E65038618B3
DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org;
	s=default; t=1701763629;
	bh=nLopImm+tL0U5tzq9B+1SFuTcC5ZMPtGC0e11Euayw4=;
	h=From:To:Subject:Date:In-Reply-To:References:From;
	b=aqeZzn+F9ycQnHxL0vMzuEmjcNpMt7FOnYFSwXkeKY5hAUj6Rq4+8Uq5oHYNj3kxT
	 59PtB/HS8HgJQ8nVtDNOQqgDHWP927D0kgLHbSCUm7OSZT/m9Nz0vbOfALolfvjeNY
	 jiqeMr5M7XkHzzek8hoEzFJmdXztjdqTbykTFBwY=
From: "tschwinge at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug target/112858] nvptx: 'unresolved symbol
 __cxa_thread_atexit_impl'
Date: Tue, 05 Dec 2023 08:07:09 +0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: changed
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: target
X-Bugzilla-Version: 14.0
X-Bugzilla-Keywords: testsuite-fail
X-Bugzilla-Severity: normal
X-Bugzilla-Who: tschwinge at gcc dot gnu.org
X-Bugzilla-Status: UNCONFIRMED
X-Bugzilla-Resolution: 
X-Bugzilla-Priority: P3
X-Bugzilla-Assigned-To: unassigned at gcc dot gnu.org
X-Bugzilla-Target-Milestone: ---
X-Bugzilla-Flags: 
X-Bugzilla-Changed-Fields: 
Message-ID: <bug-112858-4-5b8yo7rDK8@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-112858-4@http.gcc.gnu.org/bugzilla/>
References: <bug-112858-4@http.gcc.gnu.org/bugzilla/>
Content-Type: text/plain; charset="UTF-8"
Content-Transfer-Encoding: quoted-printable
X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/
Auto-Submitted: auto-generated
MIME-Version: 1.0
List-Id: <gcc-bugs.sourceware.org>

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=3D112858

--- Comment #1 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Indeed in 'build-gcc/nvptx-none/libstdc++-v3/libsupc++/atexit_thread.o' I s=
ee:

    // 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_arg=
3);
    ld.param.u32 %r34,[%value_in];
    }
    mov.u32 %r25,%r34;
    .loc 2 156 59
    bra $L8;
    $L9:
    [...]=