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 7D302385E001; Mon, 23 Mar 2020 14:21:02 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 7D302385E001
DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org;
 s=default; t=1584973262;
 bh=IR5Y0iwUE0kpwiEHEMPjkENlJPm1TUL8z798AEtvBi8=;
 h=From:To:Subject:Date:In-Reply-To:References:From;
 b=Y9q14AOxAz0yNXb/yFB+M2NrQv5xpotu74guY6/hGtw6Co/VOnM/JLFWiQa+CMvmL
 TI55vrpL0u4Uy5jLp9RRjpbZzNisVqwt+cWKQJ2S3N45zNJile24Z66IRJG3EQj9m1
 2ybXflnDb5GYyPjuCEsaZ9SjSRd2Ko5LeiPxIP/w=
From: "burnus at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug libgomp/81689] libgomp.c/target-link-1.c fails for nvptx:
 #pragma omp target link not implemented
Date: Mon, 23 Mar 2020 14:21:02 +0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: changed
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: libgomp
X-Bugzilla-Version: 8.0
X-Bugzilla-Keywords: openmp
X-Bugzilla-Severity: normal
X-Bugzilla-Who: burnus 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-81689-4-UXsn5xMmCd@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-81689-4@http.gcc.gnu.org/bugzilla/>
References: <bug-81689-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
X-BeenThere: gcc-bugs@gcc.gnu.org
X-Mailman-Version: 2.1.29
Precedence: list
List-Id: Gcc-bugs mailing list <gcc-bugs.gcc.gnu.org>
List-Unsubscribe: <http://gcc.gnu.org/mailman/options/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=unsubscribe>
List-Archive: <http://gcc.gnu.org/pipermail/gcc-bugs/>
List-Post: <mailto:gcc-bugs@gcc.gnu.org>
List-Help: <mailto:gcc-bugs-request@gcc.gnu.org?subject=help>
List-Subscribe: <http://gcc.gnu.org/mailman/listinfo/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=subscribe>
X-List-Received-Date: Mon, 23 Mar 2020 14:21:02 -0000

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

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
While the two PRs have fixed offloading for AMDGCN, nvptx still fails with =
the
same error. The generated assembler is:

 "ld.global.u64 %r27,[a$linkptr];\n"

But there is NO associated global-var declaration
similar to:

 "// BEGIN GLOBAL VAR DECL: __nvptx_uni\n"
 ".extern .shared .u32 __nvptx_uni[32];\n"

In principle, it should be generated by nvptx_output_aligned_decl =E2=80=93=
 but this
never called. It is currently called via ASM_OUTPUT_ALIGNED_DECL_COMMON and
ASM_OUTPUT_ALIGNED_DECL_LOCAL.

Without proper debugging, I have the feeling that ASM_OUTPUT_ALIGNED_BSS is
called for this =E2=80=93 but that's not defined for nvptx. Nor is
TARGET_ASM_SELECT_SECTION.=