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
