From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 48) id 070EB3857805; Fri, 15 Jul 2022 15:34:42 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 070EB3857805 From: "burnus at gcc dot gnu.org" To: gcc-bugs@gcc.gnu.org Subject: [Bug middle-end/106316] New: [OpenMP] Auto "declare target" should honor 'declare variant' kind(nohost) Date: Fri, 15 Jul 2022 15:34:42 +0000 X-Bugzilla-Reason: CC X-Bugzilla-Type: new X-Bugzilla-Watch-Reason: None X-Bugzilla-Product: gcc X-Bugzilla-Component: middle-end X-Bugzilla-Version: 13.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: bug_id short_desc product version bug_status keywords bug_severity priority component assigned_to reporter cc target_milestone Message-ID: 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 List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Fri, 15 Jul 2022 15:34:43 -0000 https://gcc.gnu.org/bugzilla/show_bug.cgi?id=3D106316 Bug ID: 106316 Summary: [OpenMP] Auto "declare target" should honor 'declare variant' kind(nohost) Product: gcc Version: 13.0 Status: UNCONFIRMED Keywords: openmp Severity: normal Priority: P3 Component: middle-end Assignee: unassigned at gcc dot gnu.org Reporter: burnus at gcc dot gnu.org CC: jakub at gcc dot gnu.org Target Milestone: --- The following fails during device lto1 with: ld: error: undefined symbol: external >>> referenced by decl-var.c:3 >>> /tmp/cc6C3J6T.o:(onhost) However, 'external' is never executed on the device. This seems to be more a 'declare variant' tagging issue than a symbol issue as it works when keeping only int onhost(void); /* Only declaration, no definition */ in the second file (and having the definition alongside 'external' in the first). =E2=86=92 omp_discover_declare_target_tgt_fn_r, omp_discover_declare_target= _fn_r functions. ---- First file. Compile with 'gcc -c first.c' ------ void external(void) { __builtin_printf("Hello World\n"); } /* See remark: int onhost(void) { external(); return 5; } */ ----------------------------------------------------- ---- Second file. Compile with 'gcc -fopenmp second.c first.o' ------ void external(void); int onhost(void) { external(); return 5; } int ondevice(void) { return 7; } #pragma omp declare variant (onhost) match (device=3D{kind(host)}) #pragma omp declare variant (ondevice) match (device=3D{kind(nohost)}) int stub (void) { return 9; } #pragma omp declare target int caller(void) { return stub(); } #pragma omp end declare target int main() { __builtin_printf("On host: %d (expected: 5)\n", caller()); #pragma omp target __builtin_printf("On device: %d (expected: 7)\n", caller()); return 0; } --------------------------------------------------------------------- * * * A variant would be: #pragma omp declare variant (ondevice) match (device=3D{kind(nohost)}) int stub (void) { external(); return 9; } where the 'stub();' is called in the target region but is still unreachable because of the 'nohost'. * * * While some generic solution would be good - i.e. if offload host is only GC= N, it would be nice if the following works - even if it will fail once -foffload=3Dnvptx is used #pragma omp declare variant (ondevice) match (construct=3D{target},device=3D{arch(gcn)}) int stub (void) { external(); return 9; } handling an explicit 'kind(host)' is very useful.=