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 83FC23858D33; Tue, 21 Nov 2023 15:57:10 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 83FC23858D33
DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org;
	s=default; t=1700582230;
	bh=FrLzyswn4jPL0ICDTJbOiVQR5J6kMR2TWidQPWe0cAE=;
	h=From:To:Subject:Date:In-Reply-To:References:From;
	b=v6ZtPApvEWa/QmbOzTLab3ocY7GKn/XW/0NgdDqBe4Sux9GNN+1TSfIE6VjHy9W8u
	 wl64499smZBauCBevZG3EgPUSzzlamPM8KQnkgDvnZx6KAYGIp8FjD1CuK8D1N5FYJ
	 7JcRvUu6uac8SLw3fYlpWB8Y4joCKIMLGsitRgdM=
From: "burnus at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug middle-end/110639] [OpenMP][5.1] Predefined firstprivate for
 pointers - attachment missing
Date: Tue, 21 Nov 2023 15:57:10 +0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: changed
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: middle-end
X-Bugzilla-Version: 14.0
X-Bugzilla-Keywords: openmp, wrong-code
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-110639-4-Oa1L1dqvgB@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-110639-4@http.gcc.gnu.org/bugzilla/>
References: <bug-110639-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=3D110639

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Testing shows that the offsets are correctly handled but that there is an
ordering problem. Example:

int
main ()
{
  int a[100] =3D {};
  int *p =3D &a[0];
  uintptr_t iptr;
  #pragma omp target map(a, iptr)
    iptr =3D (uintptr_t) p;

This will fail - as the implicitly added 'firstprivate' arrives too early at
GOMP_target_ext - before 'a' is mapped:

   map(alloc:MEM[(char *)p] [len: 0])
     map(firstprivate:p [pointer assign, bias: 0])
       map(tofrom:iptr [len: 8])
         map(tofrom:a [len: 400])

If 'a' is already present on the device (e.g. 'omp target enter data map(a)=
'),
it works.

Solution: The implicitly mapped C/C++ pointer variable 'p' must be added at=
 the
end of the clauses.=