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 150283858C39; Fri,  1 Oct 2021 10:54:55 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 150283858C39
From: "ams at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug target/102544] GCN offloading not working for
 'amdgcn-amd-amdhsa--gfx906:sramecc+:xnack-'
Date: Fri, 01 Oct 2021 10:54:54 +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: 11.0
X-Bugzilla-Keywords: openacc, openmp
X-Bugzilla-Severity: normal
X-Bugzilla-Who: ams 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-102544-4-MYa3MZxFmQ@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-102544-4@http.gcc.gnu.org/bugzilla/>
References: <bug-102544-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: <https://gcc.gnu.org/mailman/options/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=unsubscribe>
List-Archive: <https://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: <https://gcc.gnu.org/mailman/listinfo/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=subscribe>
X-List-Received-Date: Fri, 01 Oct 2021 10:54:55 -0000

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

--- Comment #3 from Andrew Stubbs <ams at gcc dot gnu.org> ---
That output shows that we have the correct libgomp and rocm is installed and
working. Libgomp initialized the GCN plugin, but did not attempt to initial=
ize
the device (the next message in the output should have been "Selected kernel
arguments memory region", or at least a GCN error message).

Instead we have a target-independent libgomp error. Presumably the kernel
metadata is malformed, somehow?

I think we need a testcase to debug this further, preferably reduced to be =
as
simple as possible.

Perhaps it would be a good idea to start with a minimal toy example and see=
 if
that works on the device.

#include <openacc.h>
#include <stdio.h>

int main ()
{
  int v =3D 1;

#pragma acc parallel copy(v)
  {
    if (acc_on_device(acc_device_host))
      v =3D -1; // error
    else {
      v =3D 2; // success
    }
  }

  printf ("v is %d\n", v);
  return v;
}=