From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-patches-return-403855-listarch-gcc-patches=gcc.gnu.org@gcc.gnu.org>
Received: (qmail 82645 invoked by alias); 24 Jul 2015 20:05:14 -0000
Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm
Precedence: bulk
List-Id: <gcc-patches.gcc.gnu.org>
List-Archive: <http://gcc.gnu.org/ml/gcc-patches/>
List-Post: <mailto:gcc-patches@gcc.gnu.org>
List-Help: <mailto:gcc-patches-help@gcc.gnu.org>
Sender: gcc-patches-owner@gcc.gnu.org
Received: (qmail 82633 invoked by uid 89); 24 Jul 2015 20:05:13 -0000
Authentication-Results: sourceware.org; auth=none
X-Virus-Found: No
X-Spam-SWARE-Status: No, score=-0.6 required=5.0 tests=AWL,BAYES_50,KAM_LAZY_DOMAIN_SECURITY,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2
X-HELO: mx1.redhat.com
Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Fri, 24 Jul 2015 20:05:06 +0000
Received: from int-mx09.intmail.prod.int.phx2.redhat.com (int-mx09.intmail.prod.int.phx2.redhat.com [10.5.11.22])	by mx1.redhat.com (Postfix) with ESMTPS id D1C1FBB997;	Fri, 24 Jul 2015 20:05:04 +0000 (UTC)
Received: from tucnak.zalov.cz (ovpn-116-43.ams2.redhat.com [10.36.116.43])	by int-mx09.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t6OK51It010385	(version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO);	Fri, 24 Jul 2015 16:05:03 -0400
Received: from tucnak.zalov.cz (localhost [127.0.0.1])	by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t6OK4xup006354;	Fri, 24 Jul 2015 22:04:59 +0200
Received: (from jakub@localhost)	by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t6OK4v1k006353;	Fri, 24 Jul 2015 22:04:57 +0200
Date: Fri, 24 Jul 2015 20:33:00 -0000
From: Jakub Jelinek <jakub@redhat.com>
To: Ilya Verbin <iverbin@gmail.com>
Cc: Thomas Schwinge <thomas@codesourcery.com>, gcc-patches@gcc.gnu.org,        Kirill Yukhin <kirill.yukhin@gmail.com>
Subject: Re: [gomp4.1] Initial support for some OpenMP 4.1 construct parsing
Message-ID: <20150724200457.GB1750@tucnak.redhat.com>
Reply-To: Jakub Jelinek <jakub@redhat.com>
References: <20150429120644.GG1751@tucnak.redhat.com> <20150609183608.GA47936@msticlxl57.ims.intel.com> <20150609202426.GG10247@tucnak.redhat.com> <20150625194529.GB33078@msticlxl57.ims.intel.com> <20150625201058.GK10247@tucnak.redhat.com> <20150717163136.GB15252@msticlxl57.ims.intel.com> <20150717164306.GT1780@tucnak.redhat.com> <20150720161422.GC1780@tucnak.redhat.com> <20150720181041.GE1780@tucnak.redhat.com> <20150722211348.GA1750@tucnak.redhat.com>
MIME-Version: 1.0
Content-Type: text/plain; charset=us-ascii
Content-Disposition: inline
In-Reply-To: <20150722211348.GA1750@tucnak.redhat.com>
User-Agent: Mutt/1.5.23 (2014-03-12)
X-IsSubscribed: yes
X-SW-Source: 2015-07/txt/msg02116.txt.bz2

On Wed, Jul 22, 2015 at 11:13:48PM +0200, Jakub Jelinek wrote:
> On Mon, Jul 20, 2015 at 08:10:41PM +0200, Jakub Jelinek wrote:
> > And here is untested incremental libgomp side of the proposed
> > GOMP_MAP_FIRSTPRIVATE_POINTER.
> 
> Actually, that seems unnecessary, for the array section maps we already
> have there a pointer, so we can easily implement that just on the
> compiler side.
> 
> Here is a WIP patch.

Another version.
What to do with zero-length array sections vs. objects is still under heated
debates, so target8.f90 keeps failing intermittently.
There is also a problem with the firstprivate implementation on #pragma omp
target for host fallback, will need to figure out something for that (the
implementation attempts to avoid double copying).  I'm considering
optimizing integral (up to bitsize of pointer)/pointer firstprivate using some new kind
GOMP_MAP_FIRSTPRIVATE_SCALAR or so, where the pointer would not be pointer
to the scalar, but the scalar itself cast to uintptr_t and then to pointer.
And then for GOMP_MAP_FIRSTPRIVATE probably even for shared space I have to
handle them (allocate using alloca, copy).

--- libgomp/testsuite/libgomp.c++/target-7.C.jj	2015-07-22 11:36:53.042867520 +0200
+++ libgomp/testsuite/libgomp.c++/target-7.C	2015-07-22 11:32:00.000000000 +0200
@@ -0,0 +1,90 @@
+extern "C" void abort ();
+
+void
+foo (int *x, int *&y, int (&z)[15])
+{
+  int a[10], b[15], err, i;
+  for (i = 0; i < 10; i++)
+    a[i] = 7 * i;
+  for (i = 0; i < 15; i++)
+    b[i] = 8 * i;
+  #pragma omp target map(to:x[5:10], y[5:10], z[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if (x[5 + i] != 20 + 4 * i
+	  || y[5 + i] != 25 + 5 * i
+	  || z[5 + i] != 30 + 6 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n, int v)
+{
+  int a[n], b[n], c[n], d[n], e[n], err, i;
+  int (*x)[n] = &c;
+  int (*y2)[n] = &d;
+  int (*&y)[n] = y2;
+  int (&z)[n] = e;
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 4 * i;
+      (*y)[i] = 5 * i;
+      z[i] = 6 * i;
+      a[i] = 7 * i;
+      b[i] = 8 * i;
+    }
+  #pragma omp target map(to:x[0][5:10], y[0][5:10], z[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 20 + 4 * i
+	  || (*y)[5 + i] != 25 + 5 * i
+	  || z[5 + i] != 30 + 6 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 9 * i;
+      (*y)[i] = 10 * i;
+      z[i] = 11 * i;
+      a[i] = 12 * i;
+      b[i] = 13 * i;
+    }
+  #pragma omp target map(to:x[0][v:v+5], y[0][v:v+5], z[v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 45 + 9 * i
+	  || (*y)[5 + i] != 50 + 10 * i
+	  || z[5 + i] != 55 + 11 * i
+	  || a[i] != 12 * i
+	  || b[5 + i] != 65 + 13 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int x[15], y2[15], z[15], *y = y2, i;
+  for (i = 0; i < 15; i++)
+    {
+      x[i] = 4 * i;
+      y[i] = 5 * i;
+      z[i] = 6 * i;
+    }
+  foo (x, y, z);
+  bar (15, 5);
+}
--- libgomp/testsuite/libgomp.c++/target-2.C.jj	2015-06-30 14:24:03.000000000 +0200
+++ libgomp/testsuite/libgomp.c++/target-2.C	2015-07-23 17:48:08.978674497 +0200
@@ -33,7 +33,8 @@ fn2 (int x, double (&dr) [1024], double
   int j;
   fn1 (hr + 2 * x, ir + 2 * x, x);
   #pragma omp target map(to: br[:x], cr[0:x], dr[x:x], er[x:x]) \
-		     map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x])
+		     map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x]) \
+		     map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (j = 0; j < x; j++)
 	s += br[j] * cr[j] + dr[x + j] + er[x + j]
--- libgomp/testsuite/libgomp.c/target-7.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-7.c	2015-07-23 17:12:33.159753962 +0200
@@ -37,63 +37,63 @@ foo (int f)
     abort ();
   #pragma omp target data device (d) map (to: h)
   {
-    #pragma omp target device (d)
+    #pragma omp target device (d) map (h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5)
       abort ();
     #pragma omp target update device (d) from (h)
   }
   #pragma omp target data if (v > 1) map (to: h)
   {
-    #pragma omp target if (v > 1)
+    #pragma omp target if (v > 1) map(h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6)
       abort ();
     #pragma omp target update if (v > 1) from (h)
   }
   #pragma omp target data device (d) if (v > 1) map (to: h)
   {
-    #pragma omp target device (d) if (v > 1)
+    #pragma omp target device (d) if (v > 1) map(h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7)
       abort ();
     #pragma omp target update device (d) if (v > 1) from (h)
   }
   #pragma omp target data if (v <= 1) map (to: h)
   {
-    #pragma omp target if (v <= 1)
+    #pragma omp target if (v <= 1) map (tofrom: h)
     if (omp_get_level () != 0 || h++ != 8)
       abort ();
     #pragma omp target update if (v <= 1) from (h)
   }
   #pragma omp target data device (d) if (v <= 1) map (to: h)
   {
-    #pragma omp target device (d) if (v <= 1)
+    #pragma omp target device (d) if (v <= 1) map (h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9)
       abort ();
     #pragma omp target update device (d) if (v <= 1) from (h)
   }
   #pragma omp target data if (0) map (to: h)
   {
-    #pragma omp target if (0)
+    #pragma omp target if (0) map (h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10)
       abort ();
     #pragma omp target update if (0) from (h)
   }
   #pragma omp target data device (d) if (0) map (to: h)
   {
-    #pragma omp target device (d) if (0)
+    #pragma omp target device (d) if (0) map (h)
     if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11)
       abort ();
     #pragma omp target update device (d) if (0) from (h)
   }
   #pragma omp target data if (1) map (to: h)
   {
-    #pragma omp target if (1)
+    #pragma omp target if (1) map (tofrom: h)
     if (omp_get_level () != 0 || h++ != 12)
       abort ();
     #pragma omp target update if (1) from (h)
   }
   #pragma omp target data device (d) if (1) map (to: h)
   {
-    #pragma omp target device (d) if (1)
+    #pragma omp target device (d) if (1) map (tofrom: h)
     if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13)
       abort ();
     #pragma omp target update device (d) if (1) from (h)
--- libgomp/testsuite/libgomp.c/target-15.c.jj	2015-07-22 11:37:11.655612690 +0200
+++ libgomp/testsuite/libgomp.c/target-15.c	2015-07-23 21:53:37.354632916 +0200
@@ -0,0 +1,74 @@
+extern void abort (void);
+
+void
+foo (int *x)
+{
+  int a[10], b[15], err, i;
+  for (i = 0; i < 10; i++)
+    a[i] = 7 * i;
+  for (i = 0; i < 15; i++)
+    b[i] = 8 * i;
+  #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if (x[5 + i] != 20 + 4 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n, int v)
+{
+  int a[n], b[n], c[n], d[n], e[n], err, i;
+  int (*x)[n] = &c;
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 4 * i;
+      a[i] = 7 * i;
+      b[i] = 8 * i;
+    }
+  #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 20 + 4 * i
+	  || a[i] != 7 * i
+	  || b[5 + i] != 40 + 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    {
+      (*x)[i] = 9 * i;
+      a[i] = 12 * i;
+      b[i] = 13 * i;
+    }
+  #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 10; i++)
+      if ((*x)[5 + i] != 45 + 9 * i
+	  || a[i] != 12 * i
+	  || b[5 + i] != 65 + 13 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int x[15], i;
+  for (i = 0; i < 15; i++)
+    x[i] = 4 * i;
+  foo (x);
+  bar (15, 5);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/target-2.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-2.c	2015-07-23 17:09:27.987350372 +0200
@@ -23,7 +23,7 @@ fn2 (int x)
   int i;
   fn1 (b, c, x);
   fn1 (e, d + x, x);
-  #pragma omp target map(to: b, c[:x], d[x:x], e)
+  #pragma omp target map(to: b, c[:x], d[x:x], e) map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
@@ -38,7 +38,7 @@ fn3 (int x)
   int i;
   fn1 (b, c, x);
   fn1 (e, d, x);
-  #pragma omp target
+  #pragma omp target map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[i];
@@ -56,7 +56,7 @@ fn4 (int x)
   #pragma omp target data map(from: b, c[:x], d[x:x], e)
     {
       #pragma omp target update to(b, c[:x], d[x:x], e)
-      #pragma omp target map(c[:x], d[x:x])
+      #pragma omp target map(c[:x], d[x:x], s)
 	#pragma omp parallel for reduction(+:s)
 	  for (i = 0; i < x; i++)
 	    {
--- libgomp/testsuite/libgomp.c/target-17.c.jj	2015-07-24 19:50:14.275109272 +0200
+++ libgomp/testsuite/libgomp.c/target-17.c	2015-07-24 19:47:57.000000000 +0200
@@ -0,0 +1,99 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[n], i, err;
+  for (i = 0; i < n; i++)
+    a[i] = 5 * i;
+  #pragma omp target map(to:a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 5 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 6 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target firstprivate (a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  int on = n;
+  #pragma omp target firstprivate (n) map(tofrom: n)
+  {
+    n++;
+  }
+  if (on != n)
+    abort ();
+  #pragma omp target map(tofrom: n) private (n)
+  {
+    n = 25;
+  }
+  if (on != n)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 8 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 9 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    a[i] += i;
+  #pragma omp target map(tofrom:a) map(from:err) private(a, i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      a[i] = 7;
+    #pragma omp parallel for reduction(|:err)
+    for (i = 0; i < n; i++)
+      if (a[i] != 7)
+        err |= 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < n; i++)
+    if (a[i] != 10 * i)
+      abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/examples-4/e.54.2.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.54.2.c	2015-07-23 16:02:02.343554209 +0200
@@ -32,7 +32,7 @@ float dotprod (float B[], float C[], int
   int i, i0;
   float sum = 0;
 
-  #pragma omp target map(to: B[0:n], C[0:n])
+  #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom: sum)
     #pragma omp teams num_teams(num_teams) thread_limit(block_threads) \
 		      reduction(+:sum)
       #pragma omp distribute
--- libgomp/testsuite/libgomp.c/examples-4/e.57.1.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.57.1.c	2015-07-23 17:37:01.880139916 +0200
@@ -10,11 +10,11 @@ int main ()
   int b = 0;
   int c, d;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
@@ -26,11 +26,11 @@ int main ()
   a += 200;
   b = 0;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
@@ -42,11 +42,11 @@ int main ()
   a += 200;
   b = 0;
 
-  #pragma omp target if(a > 200 && a < 400)
+  #pragma omp target if(a > 200 && a < 400) map(from: c)
     c = omp_is_initial_device ();
 
   #pragma omp target data map(to: b) if(a > 200 && a < 400)
-    #pragma omp target
+    #pragma omp target map(from: b, d)
       {
 	b = 100;
 	d = omp_is_initial_device ();
--- libgomp/testsuite/libgomp.c/examples-4/e.57.3.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.57.3.c	2015-07-23 16:08:48.176775074 +0200
@@ -9,7 +9,7 @@ int main ()
   int res;
   int default_device = omp_get_default_device ();
 
-  #pragma omp target
+  #pragma omp target map(from: res)
     res = omp_is_initial_device ();
 
   if (res)
@@ -17,7 +17,7 @@ int main ()
 
   omp_set_default_device (omp_get_num_devices ());
 
-  #pragma omp target
+  #pragma omp target map(from: res)
     res = omp_is_initial_device ();
 
   if (!res)
--- libgomp/testsuite/libgomp.c/examples-4/e.53.4.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.53.4.c	2015-07-23 16:00:22.468976440 +0200
@@ -41,7 +41,7 @@ float accum (int k)
   int i;
   float tmp = 0.0;
 
-  #pragma omp target
+  #pragma omp target map(tofrom:tmp)
     #pragma omp parallel for reduction(+:tmp)
       for (i = 0; i < N; i++)
 	tmp += Pfun (i, k);
--- libgomp/testsuite/libgomp.c/examples-4/e.54.4.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.54.4.c	2015-07-23 16:03:21.446427770 +0200
@@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int
   int i;
   float sum = 0;
 
-  #pragma omp target map(to: B[0:n], C[0:n])
+  #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom:sum)
     #pragma omp teams num_teams(8) thread_limit(16)
       #pragma omp distribute parallel for reduction(+:sum) \
 					  dist_schedule(static, 1024) \
--- libgomp/testsuite/libgomp.c/examples-4/e.53.5.c.jj	2015-06-17 21:00:36.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.53.5.c	2015-07-23 16:01:17.802188485 +0200
@@ -48,7 +48,7 @@ float accum ()
   int i, k;
   float tmp = 0.0;
 
-  #pragma omp target
+  #pragma omp target map(tofrom:tmp)
     #pragma omp parallel for reduction(+:tmp)
       for (i = 0; i < N; i++)
 	{
--- libgomp/testsuite/libgomp.c/examples-4/e.53.1.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.53.1.c	2015-07-23 15:59:44.430518114 +0200
@@ -20,7 +20,7 @@ int fib_wrapper (int n)
 {
   int x = 0;
 
-  #pragma omp target if(n > THRESHOLD)
+  #pragma omp target if(n > THRESHOLD) map(from:x)
     x = fib (n);
 
   return x;
--- libgomp/testsuite/libgomp.c/examples-4/e.51.3.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.51.3.c	2015-07-23 15:58:15.867779262 +0200
@@ -47,7 +47,7 @@ void gramSchmidt (int Q[][COLS], const i
       {
 	int tmp = 0;
 
-	#pragma omp target
+	#pragma omp target map(tofrom:tmp)
 	  #pragma omp parallel for reduction(+:tmp)
 	    for (i = 0; i < rows; i++)
 	      tmp += (Q[i][k] * Q[i][k]);
--- libgomp/testsuite/libgomp.c/examples-4/e.54.3.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/examples-4/e.54.3.c	2015-07-23 16:02:28.060187999 +0200
@@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int
   int i;
   float sum = 0;
 
-  #pragma omp target teams map(to: B[0:n], C[0:n])
+  #pragma omp target teams map(to: B[0:n], C[0:n]) map(tofrom: sum)
     #pragma omp distribute parallel for reduction(+:sum)
       for (i = 0; i < n; i++)
 	sum += B[i] * C[i];
--- libgomp/testsuite/libgomp.c/target-1.c.jj	2015-04-24 12:30:40.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-1.c	2015-07-23 17:08:32.474133124 +0200
@@ -34,7 +34,7 @@ fn2 (int x, int y, int z)
   fn1 (b, c, x);
   #pragma omp target data map(to: b)
   {
-    #pragma omp target map(tofrom: c)
+    #pragma omp target map(tofrom: c, s)
       #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
 	#pragma omp distribute dist_schedule(static, 4) collapse(1)
 	  for (j=0; j < x; j += y)
@@ -52,7 +52,7 @@ fn3 (int x)
   double b[1024], c[1024], s = 0;
   int i;
   fn1 (b, c, x);
-  #pragma omp target map(to: b, c)
+  #pragma omp target map(to: b, c) map(tofrom:s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	tgt (), s += b[i] * c[i];
@@ -66,7 +66,8 @@ fn4 (int x, double *p)
   int i;
   fn1 (b, c, x);
   fn1 (d + x, p + x, x);
-  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)])
+  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \
+		     map(tofrom: s)
     #pragma omp parallel for reduction(+:s)
       for (i = 0; i < x; i++)
 	s += b[i] * c[i] + d[x + i] + p[x + i];
--- libgomp/testsuite/libgomp.c/target-16.c.jj	2015-07-23 21:53:28.905753778 +0200
+++ libgomp/testsuite/libgomp.c/target-16.c	2015-07-24 12:20:32.048722516 +0200
@@ -0,0 +1,45 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+  int a[n], i, err;
+  for (i = 0; i < n; i++)
+    a[i] = 7 * i;
+  #pragma omp target firstprivate (a) map(from:err) private (i)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+}
+
+void
+bar (int n)
+{
+  int a[n], i, err;
+  #pragma omp target private (a) map(from:err)
+  {
+    #pragma omp parallel for
+    for (i = 0; i < n; i++)
+      a[i] = 7 * i;
+    err = 0;
+    #pragma omp parallel for reduction(|:err)
+    for (i = 0; i < n; i++)
+      if (a[i] != 7 * i)
+	err |= 1;
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (7);
+  bar (7);
+  return 0;
+}
--- libgomp/target.c.jj	2015-07-21 09:07:23.690851224 +0200
+++ libgomp/target.c	2015-07-22 21:12:22.438213557 +0200
@@ -142,7 +142,26 @@ resolve_device (int device_id)
 }
 
 
-/* Handle the case where splay_tree_lookup found oldn for newn.
+static inline splay_tree_key
+gomp_map_lookup (splay_tree mem_map, splay_tree_key key)
+{
+  if (key->host_start != key->host_end)
+    return splay_tree_lookup (mem_map, key);
+
+  key->host_end++;
+  splay_tree_key n = splay_tree_lookup (mem_map, key);
+  key->host_end--;
+  if (n)
+    return n;
+  key->host_start--;
+  n = splay_tree_lookup (mem_map, key);
+  key->host_start++;
+  if (n)
+    return n;
+  return splay_tree_lookup (mem_map, key);
+}
+
+/* Handle the case where gmp_map_lookup found oldn for newn.
    Helper function of gomp_map_vars.  */
 
 static inline void
@@ -204,20 +223,8 @@ gomp_map_pointer (struct target_mem_desc
     }
   /* Add bias to the pointer value.  */
   cur_node.host_start += bias;
-  cur_node.host_end = cur_node.host_start + 1;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
-  if (n == NULL)
-    {
-      /* Could be possibly zero size array section.  */
-      cur_node.host_end--;
-      n = splay_tree_lookup (mem_map, &cur_node);
-      if (n == NULL)
-	{
-	  cur_node.host_start--;
-	  n = splay_tree_lookup (mem_map, &cur_node);
-	  cur_node.host_start++;
-	}
-    }
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n == NULL)
     {
       gomp_mutex_unlock (&devicep->lock);
@@ -293,7 +300,7 @@ gomp_map_vars (struct gomp_device_descr
 	  has_firstprivate = true;
 	  continue;
 	}
-      splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+      splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
       if (n)
 	gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
 				kind & typemask);
@@ -392,7 +399,7 @@ gomp_map_vars (struct gomp_device_descr
 	      k->host_end = k->host_start + sizes[i];
 	    else
 	      k->host_end = k->host_start + sizeof (void *);
-	    splay_tree_key n = splay_tree_lookup (mem_map, k);
+	    splay_tree_key n = gomp_map_lookup (mem_map, k);
 	    if (n)
 	      gomp_map_vars_existing (devicep, n, k, &tgt->list[i],
 				      kind & typemask);
@@ -526,7 +533,8 @@ gomp_map_vars (struct gomp_device_descr
 	    }
 	  else
 	    cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
-				  + tgt->list[i].key->tgt_offset;
+				  + tgt->list[i].key->tgt_offset
+				  + tgt->list[i].offset;
 	  /* FIXME: see above FIXME comment.  */
 	  devicep->host2dev_func (devicep->target_id,
 				  (void *) (tgt->tgt_start
@@ -1289,20 +1297,8 @@ omp_target_is_present (void *ptr, size_t
   struct splay_tree_key_s cur_node;
 
   cur_node.host_start = (uintptr_t) ptr + offset;
-  cur_node.host_end = cur_node.host_start + 1;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
-  if (n == NULL)
-    {
-      /* Could be possibly zero size array section.  */
-      cur_node.host_end--;
-      n = splay_tree_lookup (mem_map, &cur_node);
-      if (n == NULL)
-	{
-	  cur_node.host_start--;
-	  n = splay_tree_lookup (mem_map, &cur_node);
-	  cur_node.host_start++;
-	}
-    }
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   int ret = n != NULL;
   gomp_mutex_unlock (&devicep->lock);
   return ret;
@@ -1524,7 +1520,7 @@ omp_target_associate_ptr (void *host_ptr
 
   cur_node.host_start = (uintptr_t) host_ptr;
   cur_node.host_end = cur_node.host_start + size;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n)
     {
       if (n->tgt->tgt_start + n->tgt_offset
@@ -1584,13 +1580,8 @@ omp_target_disassociate_ptr (void *ptr,
   int ret = EINVAL;
 
   cur_node.host_start = (uintptr_t) ptr;
-  cur_node.host_end = cur_node.host_start + 1;
-  splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
-  if (n == NULL)
-    {
-      cur_node.host_end--;
-      n = splay_tree_lookup (mem_map, &cur_node);
-    }
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
   if (n
       && n->host_start == cur_node.host_start
       && n->refcount == REFCOUNT_INFINITY
--- libgomp/libgomp.h.jj	2015-07-15 13:00:32.000000000 +0200
+++ libgomp/libgomp.h	2015-07-22 21:09:39.023307107 +0200
@@ -647,11 +647,9 @@ struct target_var_desc {
   bool copy_from;
   /* True if data always should be copied from device to host at the end.  */
   bool always_copy_from;
-  /* Used for unmapping of array sections, can be nonzero only when
-     always_copy_from is true.  */
+  /* Relative offset against key host_start.  */
   uintptr_t offset;
-  /* Used for unmapping of array sections, can be less than the size of the
-     whole object only when always_copy_from is true.  */
+  /* Actual length.  */
   uintptr_t length;
 };
 
--- include/gomp-constants.h.jj	2015-07-21 09:07:23.689851239 +0200
+++ include/gomp-constants.h	2015-07-21 15:01:05.384829637 +0200
@@ -95,7 +95,11 @@ enum gomp_map_kind
     GOMP_MAP_DELETE =			GOMP_MAP_FORCE_DEALLOC,
     /* Decrement usage count and deallocate if zero.  */
     GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_ALWAYS
-					 | GOMP_MAP_FORCE_DEALLOC)
+					 | GOMP_MAP_FORCE_DEALLOC),
+
+    /* Internal to GCC, not used in libgomp.  */
+    /* Do not map, but pointer assign a pointer instead.  */
+    GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
--- gcc/cp/parser.c.jj	2015-07-21 09:06:42.000000000 +0200
+++ gcc/cp/parser.c	2015-07-23 12:46:22.172652420 +0200
@@ -32276,27 +32276,28 @@ cp_parser_omp_target_data (cp_parser *pa
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc))
-	 {
-	 case GOMP_MAP_TO:
-	 case GOMP_MAP_ALWAYS_TO:
-	 case GOMP_MAP_FROM:
-	 case GOMP_MAP_ALWAYS_FROM:
-	 case GOMP_MAP_TOFROM:
-	 case GOMP_MAP_ALWAYS_TOFROM:
-	 case GOMP_MAP_ALLOC:
-	 case GOMP_MAP_POINTER:
-	   map_seen = 3;
-	   break;
-	 default:
-	   map_seen |= 1;
-	   error_at (OMP_CLAUSE_LOCATION (*pc),
-		     "%<#pragma omp target data%> with map-type other "
-		     "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
-		     "on %<map%> clause");
-	   *pc = OMP_CLAUSE_CHAIN (*pc);
-	   continue;
-	 }
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_ALLOC:
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target data%> with map-type other "
+		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		      "on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
 
@@ -32370,22 +32371,23 @@ cp_parser_omp_target_enter_data (cp_pars
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc))
-	 {
-	 case GOMP_MAP_TO:
-	 case GOMP_MAP_ALWAYS_TO:
-	 case GOMP_MAP_ALLOC:
-	 case GOMP_MAP_POINTER:
-	   map_seen = 3;
-	   break;
-	 default:
-	   map_seen |= 1;
-	   error_at (OMP_CLAUSE_LOCATION (*pc),
-		     "%<#pragma omp target enter data%> with map-type other "
-		     "than %<to%> or %<alloc%> on %<map%> clause");
-	   *pc = OMP_CLAUSE_CHAIN (*pc);
-	   continue;
-	 }
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_ALLOC:
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target enter data%> with map-type other "
+		      "than %<to%> or %<alloc%> on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
 
@@ -32455,24 +32457,25 @@ cp_parser_omp_target_exit_data (cp_parse
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc))
-	 {
-	 case GOMP_MAP_FROM:
-	 case GOMP_MAP_ALWAYS_FROM:
-	 case GOMP_MAP_RELEASE:
-	 case GOMP_MAP_DELETE:
-	 case GOMP_MAP_POINTER:
-	   map_seen = 3;
-	   break;
-	 default:
-	   map_seen |= 1;
-	   error_at (OMP_CLAUSE_LOCATION (*pc),
-		     "%<#pragma omp target exit data%> with map-type other "
-		     "than %<from%>, %<release%> or %<delete%> on %<map%>"
-		     " clause");
-	   *pc = OMP_CLAUSE_CHAIN (*pc);
-	   continue;
-	 }
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_RELEASE:
+	  case GOMP_MAP_DELETE:
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target exit data%> with map-type other "
+		      "than %<from%>, %<release%> or %<delete%> on %<map%>"
+		      " clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
 
@@ -32637,6 +32640,7 @@ cp_parser_omp_target (cp_parser *parser,
 	  TREE_TYPE (stmt) = void_type_node;
 	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
 	  OMP_TARGET_BODY (stmt) = body;
+	  OMP_TARGET_COMBINED (stmt) = 1;
 	  add_stmt (stmt);
 	  pc = &OMP_TARGET_CLAUSES (stmt);
 	  goto check_clauses;
@@ -32697,7 +32701,7 @@ check_clauses:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
--- gcc/cp/semantics.c.jj	2015-07-17 13:59:27.000000000 +0200
+++ gcc/cp/semantics.c	2015-07-22 13:01:26.296499686 +0200
@@ -4650,7 +4650,7 @@ handle_omp_array_sections_1 (tree c, tre
 /* Handle array sections for clause C.  */
 
 static bool
-handle_omp_array_sections (tree c)
+handle_omp_array_sections (tree c, bool is_omp)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
@@ -4828,8 +4828,9 @@ handle_omp_array_sections (tree c)
 	    return false;
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 				      OMP_CLAUSE_MAP);
-	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
-	  if (!cxx_mark_addressable (t))
+	  OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER
+					      : GOMP_MAP_POINTER);
+	  if (!is_omp && !cxx_mark_addressable (t))
 	    return false;
 	  OMP_CLAUSE_DECL (c2) = t;
 	  t = build_fold_addr_expr (first);
@@ -4847,7 +4848,8 @@ handle_omp_array_sections (tree c)
 	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
 	  OMP_CLAUSE_CHAIN (c) = c2;
 	  ptr = OMP_CLAUSE_DECL (c2);
-	  if (TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
+	  if (!is_omp
+	      && TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
 	      && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
 	    {
 	      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
@@ -5569,7 +5571,7 @@ finish_omp_clauses (tree clauses, bool a
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, allow_fields))
 		{
 		  remove = true;
 		  break;
@@ -6155,7 +6157,7 @@ finish_omp_clauses (tree clauses, bool a
 	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, allow_fields))
 		remove = true;
 	      break;
 	    }
@@ -6189,7 +6191,7 @@ finish_omp_clauses (tree clauses, bool a
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, allow_fields))
 		remove = true;
 	      else
 		{
@@ -6242,7 +6244,9 @@ finish_omp_clauses (tree clauses, bool a
 		   && !cxx_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		     && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
+		     && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+			 || (OMP_CLAUSE_MAP_KIND (c)
+			     == GOMP_MAP_FIRSTPRIVATE_POINTER)))
 		   && !type_dependent_expression_p (t)
 		   && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
 					      == REFERENCE_TYPE)
--- gcc/tree.h.jj	2015-07-16 17:56:41.000000000 +0200
+++ gcc/tree.h	2015-07-24 15:27:17.485633106 +0200
@@ -1341,6 +1341,11 @@ extern void protected_set_expr_location
 #define OMP_TEAMS_COMBINED(NODE) \
   (OMP_TEAMS_CHECK (NODE)->base.private_flag)
 
+/* True on an OMP_TARGET statement if it represents explicit
+   combined target teams, target parallel or target simd constructs.  */
+#define OMP_TARGET_COMBINED(NODE) \
+  (OMP_TARGET_CHECK (NODE)->base.private_flag)
+
 /* True if OMP_ATOMIC* is supposed to be sequentially consistent
    as opposed to relaxed.  */
 #define OMP_ATOMIC_SEQ_CST(NODE) \
@@ -1445,13 +1450,17 @@ extern void protected_set_expr_location
   ((enum gomp_map_kind) OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind)
 #define OMP_CLAUSE_SET_MAP_KIND(NODE, MAP_KIND) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \
-   = (unsigned char) (MAP_KIND))
+   = (unsigned int) (MAP_KIND))
 
 /* Nonzero if this map clause is for array (rather than pointer) based array
    section with zero bias.  Both the non-decl OMP_CLAUSE_MAP and corresponding
    OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag.  */
 #define OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.public_flag)
+/* Nonzero if the same decl appears both in OMP_CLAUSE_MAP and either
+   OMP_CLAUSE_PRIVATE or OMP_CLAUSE_FIRSTPRIVATE.  */
+#define OMP_CLAUSE_MAP_PRIVATE(NODE) \
+  TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
 
 #define OMP_CLAUSE_PROC_BIND_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)
--- gcc/gimplify.c.jj	2015-07-16 17:56:41.000000000 +0200
+++ gcc/gimplify.c	2015-07-24 17:41:57.778481242 +0200
@@ -90,6 +90,8 @@ enum gimplify_omp_var_data
   /* Flag for GOVD_LINEAR or GOVD_LASTPRIVATE: no outer reference.  */
   GOVD_LINEAR_LASTPRIVATE_NO_OUTER = 16384,
 
+  GOVD_MAP_0LEN_ARRAY = 32768,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -110,6 +112,7 @@ enum omp_region_type
   ORT_TARGET_DATA = 16,
   /* Data region with offloading.  */
   ORT_TARGET = 32,
+  ORT_COMBINED_TARGET = 33,
   /* Dummy OpenMP region, used to disable expansion of
      DECL_VALUE_EXPRs in taskloop pre body.  */
   ORT_NONE = 64
@@ -156,6 +159,9 @@ struct gimplify_omp_ctx
   enum omp_region_type region_type;
   bool combined_loop;
   bool distribute;
+  bool target_map_scalars_firstprivate;
+  bool target_map_pointers_as_0len_arrays;
+  bool target_firstprivatize_array_bases;
 };
 
 static struct gimplify_ctx *gimplify_ctxp;
@@ -2260,7 +2266,7 @@ maybe_fold_stmt (gimple_stmt_iterator *g
 {
   struct gimplify_omp_ctx *ctx;
   for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
-    if (ctx->region_type == ORT_TARGET)
+    if ((ctx->region_type & ORT_TARGET) != 0)
       return false;
   return fold_stmt (gsi);
 }
@@ -5561,8 +5567,13 @@ omp_firstprivatize_variable (struct gimp
 	  else
 	    return;
 	}
-      else if (ctx->region_type == ORT_TARGET)
-	omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+      else if ((ctx->region_type & ORT_TARGET) != 0)
+	{
+	  if (ctx->target_map_scalars_firstprivate)
+	    omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
+	  else
+	    omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+	}
       else if (ctx->region_type != ORT_WORKSHARE
 	       && ctx->region_type != ORT_SIMD
 	       && ctx->region_type != ORT_TARGET_DATA)
@@ -5648,7 +5659,7 @@ omp_add_variable (struct gimplify_omp_ct
     flags |= GOVD_SEEN;
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
-  if (n != NULL && n->value != GOVD_ALIGNED)
+  if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
     {
       /* We shouldn't be re-adding the decl with the same data
 	 sharing class.  */
@@ -5678,6 +5689,9 @@ omp_add_variable (struct gimplify_omp_ct
 	    nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
 	  else if (flags & GOVD_PRIVATE)
 	    nflags = GOVD_PRIVATE;
+	  else if ((ctx->region_type & ORT_TARGET) != 0
+		   && (flags & GOVD_FIRSTPRIVATE))
+	    nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
 	  else
 	    nflags = GOVD_FIRSTPRIVATE;
 	  nflags |= flags & GOVD_SEEN;
@@ -5746,7 +5760,7 @@ omp_notice_threadprivate_variable (struc
   struct gimplify_omp_ctx *octx;
 
   for (octx = ctx; octx; octx = octx->outer_context)
-    if (octx->region_type == ORT_TARGET)
+    if ((octx->region_type & ORT_TARGET) != 0)
       {
 	n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
 	if (n == NULL)
@@ -5810,19 +5824,66 @@ omp_notice_variable (struct gimplify_omp
     }
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
-  if (ctx->region_type == ORT_TARGET)
+  if ((ctx->region_type & ORT_TARGET) != 0)
     {
       ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
       if (n == NULL)
 	{
-	  if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+	  unsigned nflags = flags;
+	  if (ctx->target_map_pointers_as_0len_arrays
+	      || ctx->target_map_scalars_firstprivate)
+	    {
+	      bool is_declare_target = false;
+	      bool is_scalar = false;
+	      if (is_global_var (decl)
+		  && varpool_node::get_create (decl)->offloadable)
+		{
+		  struct gimplify_omp_ctx *octx;
+		  for (octx = ctx->outer_context;
+		       octx; octx = octx->outer_context)
+		    {
+		      n = splay_tree_lookup (octx->variables,
+					     (splay_tree_key)decl);
+		      if (n
+			  && (n->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED
+			  && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
+			break;
+		    }
+		  is_declare_target = octx == NULL;
+		}
+	      if (!is_declare_target && ctx->target_map_scalars_firstprivate)
+		{
+		  tree type = TREE_TYPE (decl);
+		  if (TREE_CODE (type) == REFERENCE_TYPE)
+		    type = TREE_TYPE (type);
+		  if (TREE_CODE (type) == COMPLEX_TYPE)
+		    type = TREE_TYPE (type);
+		  if (INTEGRAL_TYPE_P (type)
+		      || SCALAR_FLOAT_TYPE_P (type)
+		      || TREE_CODE (type) == POINTER_TYPE)
+		    is_scalar = true;
+		}
+	      if (is_declare_target)
+		;
+	      else if (ctx->target_map_pointers_as_0len_arrays
+		       && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+			   || (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+			       && TREE_CODE (TREE_TYPE (TREE_TYPE (decl)))
+				  == POINTER_TYPE)))
+		nflags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY;
+	      else if (is_scalar)
+		nflags |= GOVD_FIRSTPRIVATE;
+	    }
+	  if (nflags == flags
+	      && !lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
 	    {
 	      error ("%qD referenced in target region does not have "
 		     "a mappable type", decl);
-	      omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
+	      nflags |= GOVD_MAP | GOVD_EXPLICIT;
 	    }
-	  else
-	    omp_add_variable (ctx, decl, GOVD_MAP | flags);
+	  else if (nflags == flags)
+	    nflags |= GOVD_MAP;
+	  omp_add_variable (ctx, decl, nflags);
 	}
       else
 	{
@@ -6144,6 +6205,24 @@ gimplify_scan_omp_clauses (tree *list_p,
 
   ctx = new_omp_context (region_type);
   outer_ctx = ctx->outer_context;
+  if (code == OMP_TARGET && !lang_GNU_Fortran ())
+    {
+      ctx->target_map_pointers_as_0len_arrays = true;
+      /* FIXME: For Fortran we want to set this too, when
+	 the Fortran FE is updated to OpenMP 4.1.  */
+      ctx->target_map_scalars_firstprivate = true;
+    }
+  if (!lang_GNU_Fortran ())
+    switch (code)
+      {
+      case OMP_TARGET:
+      case OMP_TARGET_DATA:
+      case OMP_TARGET_ENTER_DATA:
+      case OMP_TARGET_EXIT_DATA:
+	ctx->target_firstprivatize_array_bases = true;
+      default:
+	break;
+      }
 
   while ((c = *list_p) != NULL)
     {
@@ -6290,11 +6369,18 @@ gimplify_scan_omp_clauses (tree *list_p,
 			   && ctx->region_type == ORT_WORKSHARE
 			   && octx == outer_ctx)
 		    flags = GOVD_SEEN | GOVD_SHARED;
+		  else if (octx
+			   && octx->region_type == ORT_COMBINED_TARGET)
+		    flags &= ~GOVD_LASTPRIVATE;
 		  else
 		    break;
-		  gcc_checking_assert (splay_tree_lookup (octx->variables,
-							  (splay_tree_key)
-							  decl) == NULL);
+		  splay_tree_node on
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) decl);
+		  gcc_assert (on == NULL
+			      || (octx->region_type == ORT_COMBINED_TARGET
+				  && (on->value
+				      & GOVD_DATA_SHARE_CLASS) == 0));
 		  omp_add_variable (octx, decl, flags);
 		  if (octx->outer_context == NULL)
 		    break;
@@ -6319,10 +6405,24 @@ gimplify_scan_omp_clauses (tree *list_p,
 	case OMP_CLAUSE_MAP:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (error_operand_p (decl))
+	    remove = true;
+	  switch (code)
 	    {
-	      remove = true;
+	    case OMP_TARGET:
+	      break;
+	    case OMP_TARGET_DATA:
+	    case OMP_TARGET_ENTER_DATA:
+	    case OMP_TARGET_EXIT_DATA:
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		/* For target {,enter ,exit }data only the array slice is
+		   mapped, but not the pointer to it.  */
+		remove = true;
+	      break;
+	    default:
 	      break;
 	    }
+	  if (remove)
+	    break;
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -6332,6 +6432,14 @@ gimplify_scan_omp_clauses (tree *list_p,
 	      remove = true;
 	      break;
 	    }
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+	    {
+	      OMP_CLAUSE_SIZE (c)
+		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL);
+	      omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	    }
 	  if (!DECL_P (decl))
 	    {
 	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
@@ -6643,7 +6751,10 @@ gimplify_scan_omp_clauses (tree *list_p,
 	case OMP_CLAUSE_NOGROUP:
 	case OMP_CLAUSE_THREADS:
 	case OMP_CLAUSE_SIMD:
+	  break;
+
 	case OMP_CLAUSE_DEFAULTMAP:
+	  ctx->target_map_scalars_firstprivate = false;
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -6759,6 +6870,29 @@ gimplify_adjust_omp_clauses_1 (splay_tre
     OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
     OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
+  else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
+    {
+      tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_DECL (nc) = decl;
+      if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+	  && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == POINTER_TYPE)
+	OMP_CLAUSE_DECL (clause)
+	  = build_simple_mem_ref_loc (input_location, decl);
+      OMP_CLAUSE_DECL (clause)
+	= build2 (MEM_REF, char_type_node, OMP_CLAUSE_DECL (clause),
+		  build_int_cst (build_pointer_type (char_type_node), 0));
+      OMP_CLAUSE_SIZE (clause) = size_zero_node;
+      OMP_CLAUSE_SIZE (nc) = size_zero_node;
+      OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC);
+      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      OMP_CLAUSE_CHAIN (nc) = *list_p;
+      OMP_CLAUSE_CHAIN (clause) = nc;
+      struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+      gimplify_omp_ctxp = ctx->outer_context;
+      gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0),
+		     pre_p, NULL, is_gimple_val, fb_rvalue);
+      gimplify_omp_ctxp = ctx;
+    }
   else if (code == OMP_CLAUSE_MAP)
     {
       OMP_CLAUSE_SET_MAP_KIND (clause,
@@ -6785,7 +6919,10 @@ gimplify_adjust_omp_clauses_1 (splay_tre
 				      OMP_CLAUSE_MAP);
 	  OMP_CLAUSE_DECL (nc) = decl;
 	  OMP_CLAUSE_SIZE (nc) = size_zero_node;
-	  OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+	  if (gimplify_omp_ctxp->target_firstprivatize_array_bases)
+	    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	  else
+	    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
 	  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
 	  OMP_CLAUSE_CHAIN (clause) = nc;
 	}
@@ -6910,12 +7047,14 @@ gimplify_adjust_omp_clauses (gimple_seq
 	  if (!DECL_P (decl))
 	    break;
 	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
-	  if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN)
+	  if ((ctx->region_type & ORT_TARGET) != 0
+	      && !(n->value & GOVD_SEEN)
 	      && !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS))
 	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
-		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER)
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
 	      /* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because
 		 for these, TREE_CODE (DECL_SIZE (decl)) will always be
@@ -6935,17 +7074,33 @@ gimplify_adjust_omp_clauses (gimple_seq
 		  omp_notice_variable (ctx->outer_context,
 				       OMP_CLAUSE_SIZE (c), true);
 		}
-	      tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-					  OMP_CLAUSE_MAP);
-	      OMP_CLAUSE_DECL (nc) = decl;
-	      OMP_CLAUSE_SIZE (nc) = size_zero_node;
-	      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
-	      OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
-	      OMP_CLAUSE_CHAIN (c) = nc;
-	      c = nc;
+	      if (((ctx->region_type & ORT_TARGET) != 0
+		   || !ctx->target_firstprivatize_array_bases)
+		  && ((n->value & GOVD_SEEN) == 0
+		      || (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0))
+		{
+		  tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					      OMP_CLAUSE_MAP);
+		  OMP_CLAUSE_DECL (nc) = decl;
+		  OMP_CLAUSE_SIZE (nc) = size_zero_node;
+		  if (ctx->target_firstprivatize_array_bases)
+		    OMP_CLAUSE_SET_MAP_KIND (nc,
+					     GOMP_MAP_FIRSTPRIVATE_POINTER);
+		  else
+		    OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+		  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+		  OMP_CLAUSE_CHAIN (c) = nc;
+		  c = nc;
+		}
+	    }
+	  else
+	    {
+	      if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+		OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
+	      if ((n->value & GOVD_SEEN)
+		  && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)))
+		OMP_CLAUSE_MAP_PRIVATE (c) = 1;
 	    }
-	  else if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
-	    OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
 	  break;
 
 	case OMP_CLAUSE_TO:
@@ -7888,9 +8043,11 @@ gimplify_omp_workshare (tree *expr_p, gi
     case OMP_SINGLE:
       ort = ORT_WORKSHARE;
       break;
+    case OMP_TARGET:
+      ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
+      break;
     case OACC_KERNELS:
     case OACC_PARALLEL:
-    case OMP_TARGET:
       ort = ORT_TARGET;
       break;
     case OACC_DATA:
@@ -7905,7 +8062,7 @@ gimplify_omp_workshare (tree *expr_p, gi
     }
   gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
 			     TREE_CODE (expr));
-  if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+  if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
     {
       push_gimplify_context ();
       gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
--- gcc/c/c-tree.h.jj	2015-07-01 12:50:49.000000000 +0200
+++ gcc/c/c-tree.h	2015-07-22 12:47:49.185826677 +0200
@@ -649,7 +649,7 @@ extern tree c_begin_omp_task (void);
 extern tree c_finish_omp_task (location_t, tree, tree);
 extern void c_finish_omp_cancel (location_t, tree);
 extern void c_finish_omp_cancellation_point (location_t, tree);
-extern tree c_finish_omp_clauses (tree, bool = false);
+extern tree c_finish_omp_clauses (tree, bool, bool = false);
 extern tree c_build_va_arg (location_t, tree, tree);
 extern tree c_finish_transaction (location_t, tree, int);
 extern bool c_tree_equal (tree, tree);
--- gcc/c/c-typeck.c.jj	2015-07-17 13:06:58.000000000 +0200
+++ gcc/c/c-typeck.c	2015-07-22 13:00:21.130399057 +0200
@@ -11850,7 +11850,7 @@ handle_omp_array_sections_1 (tree c, tre
 /* Handle array sections for clause C.  */
 
 static bool
-handle_omp_array_sections (tree c)
+handle_omp_array_sections (tree c, bool is_omp)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
@@ -12031,8 +12031,10 @@ handle_omp_array_sections (tree c)
 	return false;
       gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
       tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
-      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
-      if (!c_mark_addressable (t))
+      OMP_CLAUSE_SET_MAP_KIND (c2, is_omp
+				   ? GOMP_MAP_FIRSTPRIVATE_POINTER
+				   : GOMP_MAP_POINTER);
+      if (!is_omp && !c_mark_addressable (t))
 	return false;
       OMP_CLAUSE_DECL (c2) = t;
       t = build_fold_addr_expr (first);
@@ -12097,7 +12099,7 @@ c_find_omp_placeholder_r (tree *tp, int
    Remove any elements from the list that are invalid.  */
 
 tree
-c_finish_omp_clauses (tree clauses, bool declare_simd)
+c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
   bitmap_head aligned_head, map_head;
@@ -12136,7 +12138,7 @@ c_finish_omp_clauses (tree clauses, bool
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, is_omp))
 		{
 		  remove = true;
 		  break;
@@ -12496,7 +12498,7 @@ c_finish_omp_clauses (tree clauses, bool
 	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, is_omp))
 		remove = true;
 	      break;
 	    }
@@ -12519,7 +12521,7 @@ c_finish_omp_clauses (tree clauses, bool
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
-	      if (handle_omp_array_sections (c))
+	      if (handle_omp_array_sections (c, is_omp))
 		remove = true;
 	      else
 		{
@@ -12556,6 +12558,8 @@ c_finish_omp_clauses (tree clauses, bool
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		     && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 			 || (OMP_CLAUSE_MAP_KIND (c)
+			     == GOMP_MAP_FIRSTPRIVATE_POINTER)
+			 || (OMP_CLAUSE_MAP_KIND (c)
 			     == GOMP_MAP_FORCE_DEVICEPTR)))
 		   && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
 	    {
--- gcc/c/c-parser.c.jj	2015-07-21 09:06:42.000000000 +0200
+++ gcc/c/c-parser.c	2015-07-23 12:51:02.636583031 +0200
@@ -12435,7 +12435,7 @@ c_parser_oacc_all_clauses (c_parser *par
   c_parser_skip_to_pragma_eol (parser);
 
   if (finish_p)
-    return c_finish_omp_clauses (clauses);
+    return c_finish_omp_clauses (clauses, false);
 
   return clauses;
 }
@@ -12720,8 +12720,8 @@ c_parser_omp_all_clauses (c_parser *pars
   if (finish_p)
     {
       if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0)
-	return c_finish_omp_clauses (clauses, true);
-      return c_finish_omp_clauses (clauses);
+	return c_finish_omp_clauses (clauses, true, true);
+      return c_finish_omp_clauses (clauses, true);
     }
 
   return clauses;
@@ -12755,7 +12755,7 @@ c_parser_oacc_cache (location_t loc, c_p
   tree stmt, clauses;
 
   clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
-  clauses = c_finish_omp_clauses (clauses);
+  clauses = c_finish_omp_clauses (clauses, false);
 
   c_parser_skip_to_pragma_eol (parser);
 
@@ -13902,7 +13902,7 @@ omp_split_clauses (location_t loc, enum
   c_omp_split_clauses (loc, code, mask, clauses, cclauses);
   for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
     if (cclauses[i])
-      cclauses[i] = c_finish_omp_clauses (cclauses[i]);
+      cclauses[i] = c_finish_omp_clauses (cclauses[i], true);
 }
 
 /* OpenMP 4.0:
@@ -14668,9 +14668,10 @@ c_parser_omp_target_data (location_t loc
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -14800,9 +14801,10 @@ c_parser_omp_target_enter_data (location
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -14885,9 +14887,10 @@ c_parser_omp_target_exit_data (location_
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_RELEASE:
 	  case GOMP_MAP_DELETE:
-	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -15016,6 +15019,7 @@ c_parser_omp_target (c_parser *parser, e
 	  TREE_TYPE (stmt) = void_type_node;
 	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
 	  OMP_TARGET_BODY (stmt) = block;
+	  OMP_TARGET_COMBINED (stmt) = 1;
 	  add_stmt (stmt);
 	  pc = &OMP_TARGET_CLAUSES (stmt);
 	  goto check_clauses;
@@ -15078,7 +15082,7 @@ check_clauses:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_ALLOC:
-	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -16379,7 +16383,7 @@ c_parser_cilk_for (c_parser *parser, tre
   tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE);
   OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR;
   OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain;
-  clauses = c_finish_omp_clauses (clauses);
+  clauses = c_finish_omp_clauses (clauses, false);
 
   tree block = c_begin_compound_stmt (true);
   tree sb = push_stmt_list ();
@@ -16444,7 +16448,7 @@ c_parser_cilk_for (c_parser *parser, tre
       OMP_CLAUSE_OPERAND (c, 0)
 	= cilk_for_number_of_iterations (omp_for);
       OMP_CLAUSE_CHAIN (c) = clauses;
-      OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c);
+      OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true);
       add_stmt (omp_par);
     }
 
--- gcc/tree-core.h.jj	2015-07-17 09:30:44.000000000 +0200
+++ gcc/tree-core.h	2015-07-21 16:28:48.524156167 +0200
@@ -1354,7 +1354,7 @@ struct GTY(()) tree_omp_clause {
     enum omp_clause_schedule_kind  schedule_kind;
     enum omp_clause_depend_kind    depend_kind;
     /* See include/gomp-constants.h for enum gomp_map_kind's values.  */
-    unsigned char		   map_kind;
+    unsigned int		   map_kind;
     enum omp_clause_proc_bind_kind proc_bind_kind;
     enum tree_code                 reduction_code;
     enum omp_clause_linear_kind    linear_kind;
--- gcc/omp-low.c.jj	2015-07-21 09:07:23.000000000 +0200
+++ gcc/omp-low.c	2015-07-24 18:12:01.474522499 +0200
@@ -1071,24 +1071,35 @@ lookup_field (tree var, omp_context *ctx
 }
 
 static inline tree
-lookup_sfield (tree var, omp_context *ctx)
+lookup_sfield (splay_tree_key key, omp_context *ctx)
 {
   splay_tree_node n;
   n = splay_tree_lookup (ctx->sfield_map
-			 ? ctx->sfield_map : ctx->field_map,
-			 (splay_tree_key) var);
+			 ? ctx->sfield_map : ctx->field_map, key);
   return (tree) n->value;
 }
 
 static inline tree
-maybe_lookup_field (tree var, omp_context *ctx)
+lookup_sfield (tree var, omp_context *ctx)
+{
+  return lookup_sfield ((splay_tree_key) var, ctx);
+}
+
+static inline tree
+maybe_lookup_field (splay_tree_key key, omp_context *ctx)
 {
   splay_tree_node n;
-  n = splay_tree_lookup (ctx->field_map, (splay_tree_key) var);
+  n = splay_tree_lookup (ctx->field_map, key);
   return n ? (tree) n->value : NULL_TREE;
 }
 
 static inline tree
+maybe_lookup_field (tree var, omp_context *ctx)
+{
+  return maybe_lookup_field ((splay_tree_key) var, ctx);
+}
+
+static inline tree
 lookup_oacc_reduction (const char *id, omp_context *ctx)
 {
   splay_tree_node n;
@@ -1359,12 +1370,18 @@ build_outer_var_ref (tree var, omp_conte
 /* Build tree nodes to access the field for VAR on the sender side.  */
 
 static tree
-build_sender_ref (tree var, omp_context *ctx)
+build_sender_ref (splay_tree_key key, omp_context *ctx)
 {
-  tree field = lookup_sfield (var, ctx);
+  tree field = lookup_sfield (key, ctx);
   return omp_build_component_ref (ctx->sender_decl, field);
 }
 
+static tree
+build_sender_ref (tree var, omp_context *ctx)
+{
+  return build_sender_ref ((splay_tree_key) var, ctx);
+}
+
 /* Add a new field for VAR inside the structure CTX->SENDER_DECL.  */
 
 static void
@@ -1908,6 +1925,10 @@ scan_sharing_clauses (tree clauses, omp_
 	case OMP_CLAUSE_LINEAR:
 	  decl = OMP_CLAUSE_DECL (c);
 	do_private:
+	  if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+	       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+	      && is_gimple_omp_offloaded (ctx->stmt))
+	    install_var_field (decl, !is_reference (decl), 3, ctx);
 	  if (is_variable_sized (decl))
 	    {
 	      if (is_task_ctx (ctx))
@@ -1930,10 +1951,6 @@ scan_sharing_clauses (tree clauses, omp_
 	      else if (!global)
 		install_var_field (decl, by_ref, 3, ctx);
 	    }
-	  else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
-		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
-		   && is_gimple_omp_offloaded (ctx->stmt))
-	    install_var_field (decl, !is_reference (decl), 3, ctx);
 	  install_var_local (decl, ctx);
 	  if (is_gimple_omp_oacc (ctx->stmt)
 	      && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
@@ -2025,6 +2042,21 @@ scan_sharing_clauses (tree clauses, omp_
 		  && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
 		break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+	    {
+	      if (DECL_SIZE (decl)
+		  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+		{
+		  tree decl2 = DECL_VALUE_EXPR (decl);
+		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+		  decl2 = TREE_OPERAND (decl2, 0);
+		  gcc_assert (DECL_P (decl2));
+		  install_var_local (decl2, ctx);
+		}
+	      install_var_local (decl, ctx);
+	      break;
+	    }
 	  if (DECL_P (decl))
 	    {
 	      if (DECL_SIZE (decl)
@@ -2034,7 +2066,11 @@ scan_sharing_clauses (tree clauses, omp_
 		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
 		  decl2 = TREE_OPERAND (decl2, 0);
 		  gcc_assert (DECL_P (decl2));
-		  install_var_field (decl2, true, 3, ctx);
+		  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		      && OMP_CLAUSE_MAP_PRIVATE (c))
+		    install_var_field (decl2, true, 11, ctx);
+		  else
+		    install_var_field (decl2, true, 3, ctx);
 		  install_var_local (decl2, ctx);
 		  install_var_local (decl, ctx);
 		}
@@ -2045,6 +2081,9 @@ scan_sharing_clauses (tree clauses, omp_
 		      && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
 		      && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 		    install_var_field (decl, true, 7, ctx);
+		  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			   && OMP_CLAUSE_MAP_PRIVATE (c))
+		    install_var_field (decl, true, 11, ctx);
 		  else
 		    install_var_field (decl, true, 3, ctx);
 		  if (is_gimple_omp_offloaded (ctx->stmt))
@@ -2151,7 +2190,19 @@ scan_sharing_clauses (tree clauses, omp_
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (is_variable_sized (decl))
-	    install_var_local (decl, ctx);
+	    {
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+		  && is_gimple_omp_offloaded (ctx->stmt))
+		{
+		  tree decl2 = DECL_VALUE_EXPR (decl);
+		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+		  decl2 = TREE_OPERAND (decl2, 0);
+		  gcc_assert (DECL_P (decl2));
+		  install_var_local (decl2, ctx);
+		  fixup_remapped_decl (decl2, ctx, false);
+		}
+	      install_var_local (decl, ctx);
+	    }
 	  fixup_remapped_decl (decl, ctx,
 			       OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
 			       && OMP_CLAUSE_PRIVATE_DEBUG (c));
@@ -2201,7 +2252,8 @@ scan_sharing_clauses (tree clauses, omp_
 	    break;
 	  if (DECL_P (decl))
 	    {
-	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+		   || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 		  && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
 		  && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
 		{
@@ -3924,11 +3976,8 @@ handle_simd_reference (location_t loc, t
   tree z = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_vard)));
   if (TREE_CONSTANT (z))
     {
-      const char *name = NULL;
-      if (DECL_NAME (new_vard))
-	name = IDENTIFIER_POINTER (DECL_NAME (new_vard));
-
-      z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)), name);
+      z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)),
+			      get_name (new_vard));
       gimple_add_tmp_var (z);
       TREE_ADDRESSABLE (z) = 1;
       z = build_fold_addr_expr_loc (loc, z);
@@ -4127,9 +4176,7 @@ lower_rec_input_clauses (tree clauses, g
 	      tree type = TREE_TYPE (d);
 	      gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
 	      tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
-	      const char *name = NULL;
-	      if (DECL_NAME (orig_var))
-		name = IDENTIFIER_POINTER (DECL_NAME (orig_var));
+	      const char *name = get_name (orig_var);
 	      if (TREE_CONSTANT (v))
 		{
 		  x = create_tmp_var_raw (type, name);
@@ -4139,7 +4186,8 @@ lower_rec_input_clauses (tree clauses, g
 		}
 	      else
 		{
-		  tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
+		  tree atmp
+		    = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
 		  tree t = maybe_lookup_decl (v, ctx);
 		  if (t)
 		    v = t;
@@ -4152,7 +4200,8 @@ lower_rec_input_clauses (tree clauses, g
 		  t = fold_build2_loc (clause_loc, MULT_EXPR,
 				       TREE_TYPE (v), t,
 				       TYPE_SIZE_UNIT (TREE_TYPE (type)));
-		  x = build_call_expr_loc (clause_loc, atmp, 1, t);
+		  tree al = size_int (TYPE_ALIGN (TREE_TYPE (type)));
+		  x = build_call_expr_loc (clause_loc, atmp, 2, t, al);
 		}
 
 	      tree ptype = build_pointer_type (TREE_TYPE (type));
@@ -4362,8 +4411,9 @@ lower_rec_input_clauses (tree clauses, g
 		  x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
 
 		  /* void *tmp = __builtin_alloca */
-		  atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
-		  stmt = gimple_build_call (atmp, 1, x);
+		  atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		  stmt = gimple_build_call (atmp, 2, x,
+					    size_int (DECL_ALIGN (var)));
 		  tmp = create_tmp_var_raw (ptr_type_node);
 		  gimple_add_tmp_var (tmp);
 		  gimple_call_set_lhs (stmt, tmp);
@@ -4400,12 +4450,8 @@ lower_rec_input_clauses (tree clauses, g
 		    x = NULL_TREE;
 		  else
 		    {
-		      const char *name = NULL;
-		      if (DECL_NAME (var))
-			name = IDENTIFIER_POINTER (DECL_NAME (new_var));
-
 		      x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
-					      name);
+					      get_name (var));
 		      gimple_add_tmp_var (x);
 		      TREE_ADDRESSABLE (x) = 1;
 		      x = build_fold_addr_expr_loc (clause_loc, x);
@@ -4413,8 +4459,11 @@ lower_rec_input_clauses (tree clauses, g
 		}
 	      else
 		{
-		  tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
-		  x = build_call_expr_loc (clause_loc, atmp, 1, x);
+		  tree atmp
+		    = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		  tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+		  tree al = size_int (TYPE_ALIGN (rtype));
+		  x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
 		}
 
 	      if (x)
@@ -5489,11 +5538,7 @@ lower_send_clauses (tree clauses, gimple
 	  /* Handle taskloop firstprivate/lastprivate, where the
 	     lastprivate on GIMPLE_OMP_TASK is represented as
 	     OMP_CLAUSE_SHARED_FIRSTPRIVATE.  */
-	  tree f
-	    = (tree)
-	      splay_tree_lookup (ctx->sfield_map
-				 ? ctx->sfield_map : ctx->field_map,
-				 (splay_tree_key) &DECL_UID (val))->value;
+	  tree f = lookup_sfield ((splay_tree_key) &DECL_UID (val), ctx);
 	  x = omp_build_component_ref (ctx->sender_decl, f);
 	  if (use_pointer_for_field (val, ctx))
 	    var = build_fold_addr_expr (var);
@@ -12883,6 +12928,7 @@ lower_omp_target (gimple_stmt_iterator *
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    break;
 	  case GOMP_MAP_FORCE_ALLOC:
 	  case GOMP_MAP_FORCE_TO:
@@ -12918,6 +12964,28 @@ lower_omp_target (gimple_stmt_iterator *
 	    var = var2;
 	  }
 
+	if (offloaded
+	    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+	  {
+	    if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+	      {
+		tree type = build_pointer_type (TREE_TYPE (var));
+		tree new_var = lookup_decl (var, ctx);
+		x = create_tmp_var_raw (type, get_name (new_var));
+		gimple_add_tmp_var (x);
+		x = build_simple_mem_ref (x);
+		SET_DECL_VALUE_EXPR (new_var, x);
+		DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	      }
+	    continue;
+	  }
+
+	if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c))
+	  {
+	    map_cnt++;
+	    continue;
+	  }
+
 	if (!maybe_lookup_field (var, ctx))
 	  continue;
 
@@ -12925,6 +12993,7 @@ lower_omp_target (gimple_stmt_iterator *
 	  {
 	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
+
 	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 		&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
 		&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
@@ -12942,8 +13011,36 @@ lower_omp_target (gimple_stmt_iterator *
 	if (!is_reference (var)
 	    && !is_gimple_reg_type (TREE_TYPE (var)))
 	  {
-	    x = build_receiver_ref (var, true, ctx);
 	    tree new_var = lookup_decl (var, ctx);
+	    if (is_variable_sized (var))
+	      {
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_pvar = lookup_decl (pvar, ctx);
+		x = build_fold_indirect_ref (new_pvar);
+		TREE_THIS_NOTRAP (x) = 1;
+	      }
+	    else
+	      x = build_receiver_ref (var, true, ctx);
+	    SET_DECL_VALUE_EXPR (new_var, x);
+	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+	  }
+	break;
+
+      case OMP_CLAUSE_PRIVATE:
+	var = OMP_CLAUSE_DECL (c);
+	if (is_variable_sized (var))
+	  {
+	    tree new_var = lookup_decl (var, ctx);
+	    tree pvar = DECL_VALUE_EXPR (var);
+	    gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+	    pvar = TREE_OPERAND (pvar, 0);
+	    gcc_assert (DECL_P (pvar));
+	    tree new_pvar = lookup_decl (pvar, ctx);
+	    x = build_fold_indirect_ref (new_pvar);
+	    TREE_THIS_NOTRAP (x) = 1;
 	    SET_DECL_VALUE_EXPR (new_var, x);
 	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
 	  }
@@ -13044,6 +13141,10 @@ lower_omp_target (gimple_stmt_iterator *
 	      }
 	    else
 	      {
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_KIND (c)
+		       == GOMP_MAP_FIRSTPRIVATE_POINTER)
+		  break;
 		if (DECL_SIZE (ovar)
 		    && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
 		  {
@@ -13053,7 +13154,14 @@ lower_omp_target (gimple_stmt_iterator *
 		    gcc_assert (DECL_P (ovar2));
 		    ovar = ovar2;
 		  }
-		if (!maybe_lookup_field (ovar, ctx))
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_PRIVATE (c))
+		  {
+		    if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar),
+					     ctx))
+		      continue;
+		  }
+		else if (!maybe_lookup_field (ovar, ctx))
 		  continue;
 	      }
 
@@ -13063,7 +13171,12 @@ lower_omp_target (gimple_stmt_iterator *
 	    if (nc)
 	      {
 		var = lookup_decl_in_outer_ctx (ovar, ctx);
-		x = build_sender_ref (ovar, ctx);
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		    && OMP_CLAUSE_MAP_PRIVATE (c))
+		  x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar),
+					ctx);
+		else
+		  x = build_sender_ref (ovar, ctx);
 		if (maybe_lookup_oacc_reduction (var, ctx))
 		  {
 		    gcc_checking_assert (offloaded
@@ -13101,7 +13214,7 @@ lower_omp_target (gimple_stmt_iterator *
 			 || map_kind == GOMP_MAP_FORCE_DEVICEPTR)
 			&& !TYPE_READONLY (TREE_TYPE (var)))
 		      {
-			x = build_sender_ref (ovar, ctx);
+			x = unshare_expr (x);
 			x = build_simple_mem_ref (x);
 			gimplify_assign (var, x, &olist);
 		      }
@@ -13239,6 +13352,7 @@ lower_omp_target (gimple_stmt_iterator *
 
   if (offloaded)
     {
+      tree prev = NULL_TREE;
       for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
 	  {
@@ -13257,6 +13371,18 @@ lower_omp_target (gimple_stmt_iterator *
 		gimple_seq_add_stmt (&new_body,
 				     gimple_build_assign (new_var, x));
 	      }
+            else if (is_variable_sized (var))
+	      {
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_var = lookup_decl (pvar, ctx);
+		tree x = build_receiver_ref (var, false, ctx);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_var, x));
+	      }
 	    break;
 	  case OMP_CLAUSE_PRIVATE:
 	    var = OMP_CLAUSE_DECL (c);
@@ -13267,20 +13393,19 @@ lower_omp_target (gimple_stmt_iterator *
 		tree x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
 		if (TREE_CONSTANT (x))
 		  {
-		    const char *name = NULL;
-		    if (DECL_NAME (var))
-		      name = IDENTIFIER_POINTER (DECL_NAME (new_var));
-
 		    x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
-					    name);
+					    get_name (var));
 		    gimple_add_tmp_var (x);
 		    TREE_ADDRESSABLE (x) = 1;
 		    x = build_fold_addr_expr_loc (clause_loc, x);
 		  }
 		else
 		  {
-		    tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
-		    x = build_call_expr_loc (clause_loc, atmp, 1, x);
+		    tree atmp
+		      = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		    tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+		    tree al = size_int (TYPE_ALIGN (rtype));
+		    x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
 		  }
 
 		x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
@@ -13290,6 +13415,110 @@ lower_omp_target (gimple_stmt_iterator *
 	      }
 	    break;
 	  }
+      /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass,
+	 so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
+	 are already handled.  */
+      for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+	switch (OMP_CLAUSE_CODE (c))
+	  {
+	    tree var;
+	  default:
+	    break;
+	  case OMP_CLAUSE_MAP:
+	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+	      {
+		location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+		gcc_assert (prev);
+		var = OMP_CLAUSE_DECL (c);
+		if (DECL_SIZE (var)
+		    && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+		  {
+		    tree var2 = DECL_VALUE_EXPR (var);
+		    gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+		    var2 = TREE_OPERAND (var2, 0);
+		    gcc_assert (DECL_P (var2));
+		    var = var2;
+		  }
+		tree new_var = lookup_decl (var, ctx), x;
+		tree type = TREE_TYPE (new_var);
+		bool is_ref = is_reference (var);
+		bool ref_to_array = false;
+		if (is_ref)
+		  {
+		    type = TREE_TYPE (type);
+		    if (TREE_CODE (type) == ARRAY_TYPE)
+		      {
+			type = build_pointer_type (type);
+			ref_to_array = true;
+		      }
+		  }
+		else if (TREE_CODE (type) == ARRAY_TYPE)
+		  {
+		    tree decl2 = DECL_VALUE_EXPR (new_var);
+		    gcc_assert (TREE_CODE (decl2) == MEM_REF);
+		    decl2 = TREE_OPERAND (decl2, 0);
+		    gcc_assert (DECL_P (decl2));
+		    new_var = decl2;
+		    type = TREE_TYPE (new_var);
+		  }
+		x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx);
+		x = fold_convert_loc (clause_loc, type, x);
+		if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
+		  {
+		    tree bias = OMP_CLAUSE_SIZE (c);
+		    if (DECL_P (bias))
+		      bias = lookup_decl (bias, ctx);
+		    bias = fold_convert_loc (clause_loc, sizetype, bias);
+		    bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype,
+					    bias);
+		    x = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR,
+					 TREE_TYPE (x), x, bias);
+		  }
+		if (ref_to_array)
+		  x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		if (is_ref && !ref_to_array)
+		  {
+		    tree t = create_tmp_var_raw (type, get_name (var));
+		    gimple_add_tmp_var (t);
+		    TREE_ADDRESSABLE (t) = 1;
+		    gimple_seq_add_stmt (&new_body,
+					 gimple_build_assign (t, x));
+		    x = build_fold_addr_expr_loc (clause_loc, t);
+		  }
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_var, x));
+		prev = NULL_TREE;
+	      }
+	    else if (OMP_CLAUSE_CHAIN (c)
+		     && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+			== OMP_CLAUSE_MAP
+		     && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			== GOMP_MAP_FIRSTPRIVATE_POINTER)
+	      prev = c;
+	    break;
+	  case OMP_CLAUSE_PRIVATE:
+	    var = OMP_CLAUSE_DECL (c);
+	    if (is_variable_sized (var))
+	      {
+		location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+		tree new_var = lookup_decl (var, ctx);
+		tree pvar = DECL_VALUE_EXPR (var);
+		gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+		pvar = TREE_OPERAND (pvar, 0);
+		gcc_assert (DECL_P (pvar));
+		tree new_pvar = lookup_decl (pvar, ctx);
+		tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+		tree al = size_int (DECL_ALIGN (var));
+		tree x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
+		x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
+		x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+		gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+		gimple_seq_add_stmt (&new_body,
+				     gimple_build_assign (new_pvar, x));
+	      }
+	    break;
+	  }
       gimple_seq_add_seq (&new_body, tgt_body);
       new_body = maybe_catch_exception (new_body);
     }
--- gcc/tree-pretty-print.c.jj	2015-07-21 09:06:42.000000000 +0200
+++ gcc/tree-pretty-print.c	2015-07-22 13:53:51.406065024 +0200
@@ -639,6 +639,9 @@ dump_omp_clause (pretty_printer *pp, tre
 	case GOMP_MAP_RELEASE:
 	  pp_string (pp, "release");
 	  break;
+	case GOMP_MAP_FIRSTPRIVATE_POINTER:
+	  pp_string (pp, "firstprivate");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -649,7 +652,9 @@ dump_omp_clause (pretty_printer *pp, tre
       if (OMP_CLAUSE_SIZE (clause))
 	{
 	  if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
-	      && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER)
+	      && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER
+		  || OMP_CLAUSE_MAP_KIND (clause)
+		     == GOMP_MAP_FIRSTPRIVATE_POINTER))
 	    pp_string (pp, " [pointer assign, bias: ");
 	  else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
 		   && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET)


	Jakub