Merge branch 'master' into devel/sphinx

11 files changed, 521 insertions, 9 deletions

diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
index 7151bd1..e199a18 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
@@ -17,6 +17,7 @@
 
 #include <omp.h>
 #include <stdlib.h>
+#include <string.h>
 
 int
 main ()
@@ -25,21 +26,28 @@ main ()
   int chunk_size;
   omp_get_schedule(&kind, &chunk_size);
 
-  if (omp_get_max_teams () != 42
-      || !omp_get_dynamic ()
-      || kind != 3 || chunk_size != 4
-      || omp_get_teams_thread_limit () != 44
-      || omp_get_thread_limit () != 45
-      || omp_get_max_threads () != 46
-      || omp_get_proc_bind () != omp_proc_bind_spread
-      || omp_get_max_active_levels () != 47)
+  if ((!getenv ("OMP_NUM_TEAMS") && omp_get_max_teams () != 42)
+      || (!getenv ("OMP_DYNAMIC") && !omp_get_dynamic ())
+      || (!getenv ("OMP_SCHEDULE") && (kind != 3 || chunk_size != 4))
+      || (!getenv ("OMP_TEAMS_THREAD_LIMIT") && omp_get_teams_thread_limit () != 44)
+      || (!getenv ("OMP_THREAD_LIMIT") && omp_get_thread_limit () != 45)
+      || (!getenv ("OMP_NUM_THREADS") && omp_get_max_threads () != 46)
+      || (!getenv ("OMP_PROC_BIND") && omp_get_proc_bind () != omp_proc_bind_spread)
+      || (!getenv ("OMP_MAX_ACTIVE_LEVELS") && omp_get_max_active_levels () != 47))
     abort ();
 
   int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
-  for (int i=0; i < num_devices; i++)
+  for (int i = 0; i < num_devices; i++)
+    {
+      char name[sizeof ("OMP_NUM_TEAMS_DEV_1")];
+      strcpy (name, "OMP_NUM_TEAMS_DEV_1");
+      name[sizeof ("OMP_NUM_TEAMS_DEV_1") - 2] = '0' + i;
+      if (getenv (name))
+	continue;
     #pragma omp target device (i)
       if (omp_get_max_teams () != 43)
 	abort ();
+    }
 
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
index 6ed5a5f..5883eff 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
@@ -1,4 +1,5 @@
 /* { dg-additional-options "-flto" } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources requires-4-aux.c } */
 
 /* Check no diagnostic by device-compiler's or host compiler's lto1.
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
index 7fe0c73..d43d78d 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources requires-5-aux.c } */
 
 /* Depending on offload device capabilities, it may print something like the
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
index b00c745..a25b4d2 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
@@ -1,3 +1,5 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+
 #pragma omp requires unified_shared_memory, unified_address, reverse_offload
 
 /* The requires line is not active as there is none of:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
index 976e129..52d828c 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
@@ -1,4 +1,5 @@
 /* { dg-do run }  */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources reverse-offload-1-aux.c } */
 
 /* Check that reverse offload works in particular:
diff --git a/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
new file mode 100644
index 0000000..fbfeae1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
@@ -0,0 +1,15 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-march=sm_30 -foffload=-mptx=_" } */
+
+#pragma omp requires reverse_offload
+
+int
+main ()
+{
+  #pragma omp target
+    {
+    }
+  return 0;
+}
+
+/* { dg-warning "'omp requires reverse_offload' requires at least 'sm_35' for '-foffload-options=nvptx-none=-march=' - disabling offload-code generation for this device type" "" { target *-*-* } 0 } */
diff --git a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90 b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
index 7cfb8b6..de68011 100644
--- a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
@@ -1,4 +1,5 @@
 ! { dg-do run }
+! { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } }
 ! { dg-additional-sources reverse-offload-1-aux.f90 }
 
 ! Check that reverse offload works in particular:
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
new file mode 100644
index 0000000..dacbb52
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
@@ -0,0 +1,101 @@
+#include <cassert>
+
+/* Test attach/detach operation with pointers and references to structs.  */
+
+typedef struct mystruct {
+  int *a;
+  int b;
+  int *c;
+  int d;
+  int *e;
+} mystruct;
+
+void str (void)
+{
+  int a[10], c[10], e[10];
+  mystruct m = { .a = a, .c = c, .e = e };
+  a[0] = 5;
+  c[0] = 7;
+  e[0] = 9;
+  #pragma acc parallel copy(m.a[0:10], m.b, m.c[0:10], m.d, m.e[0:10])
+  {
+    m.a[0] = m.c[0] + m.e[0];
+  }
+  assert (m.a[0] == 7 + 9);
+}
+
+void strp (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct *m = new mystruct;
+  m->a = a;
+  m->c = c;
+  m->e = e;
+  a[0] = 6;
+  c[0] = 8;
+  e[0] = 10;
+  #pragma acc parallel copy(m->a[0:10], m->b, m->c[0:10], m->d, m->e[0:10])
+  {
+    m->a[0] = m->c[0] + m->e[0];
+  }
+  assert (m->a[0] == 8 + 10);
+  delete m;
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strr (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct m;
+  mystruct &n = m;
+  n.a = a;
+  n.c = c;
+  n.e = e;
+  a[0] = 7;
+  c[0] = 9;
+  e[0] = 11;
+  #pragma acc parallel copy(n.a[0:10], n.b, n.c[0:10], n.d, n.e[0:10])
+  {
+    n.a[0] = n.c[0] + n.e[0];
+  }
+  assert (n.a[0] == 9 + 11);
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strrp (void)
+{
+  int a[10], c[10], e[10];
+  mystruct *m = new mystruct;
+  mystruct *&n = m;
+  n->a = a;
+  n->b = 3;
+  n->c = c;
+  n->d = 5;
+  n->e = e;
+  a[0] = 8;
+  c[0] = 10;
+  e[0] = 12;
+  #pragma acc parallel copy(n->a[0:10], n->c[0:10], n->e[0:10])
+  {
+    n->a[0] = n->c[0] + n->e[0];
+  }
+  assert (n->a[0] == 10 + 12);
+  delete m;
+}
+
+int main (int argc, char *argv[])
+{
+  str ();
+  strp ();
+  strr ();
+  strrp ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
new file mode 100644
index 0000000..27fe1a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
@@ -0,0 +1,68 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+   non-zero.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+	m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
new file mode 100644
index 0000000..a7308e8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
@@ -0,0 +1,231 @@
+#include <stdlib.h>
+
+/* Test mapping chained indirect struct accesses, mixed in different ways.  */
+
+typedef struct {
+  int *a;
+  int b;
+  int *c;
+} str1;
+
+typedef struct {
+  int d;
+  int *e;
+  str1 *f;
+} str2;
+
+typedef struct {
+  int g;
+  int h;
+  str2 *s2;
+} str3;
+
+typedef struct {
+  str3 m;
+  str3 n;
+} str4;
+
+void
+zero_arrays (str4 *s, int N)
+{
+  for (int i = 0; i < N; i++)
+    {
+      s->m.s2->e[i] = 0;
+      s->m.s2->f->a[i] = 0;
+      s->m.s2->f->c[i] = 0;
+      s->n.s2->e[i] = 0;
+      s->n.s2->f->a[i] = 0;
+      s->n.s2->f->c[i] = 0;
+    }
+}
+
+void
+alloc_s2 (str2 **s, int N)
+{
+  (*s) = (str2 *) malloc (sizeof (str2));
+  (*s)->f = (str1 *) malloc (sizeof (str1));
+  (*s)->e = (int *) malloc (sizeof (int) * N);
+  (*s)->f->a = (int *) malloc (sizeof (int) * N);
+  (*s)->f->c = (int *) malloc (sizeof (int) * N);
+}
+
+int main (int argc, char* argv[])
+{
+  const int N = 1024;
+  str4 p, *q;
+  int i;
+
+  alloc_s2 (&p.m.s2, N);
+  alloc_s2 (&p.n.s2, N);
+  q = (str4 *) malloc (sizeof (str4));
+  alloc_s2 (&q->m.s2, N);
+  alloc_s2 (&q->n.s2, N);
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc parallel loop copy(p.m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	p.m.s2->e[j]++;
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \
+			  copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99
+	|| p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \
+		       copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \
+		      copyout(p.n.s2->e[:N])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99
+	|| p.n.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc parallel loop copy(q->m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	q->m.s2->e[j]++;
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \
+			  copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99
+	|| q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \
+		       copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \
+		      copyout(q->n.s2->e[:N])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99
+	|| q->n.s2->e[i] != 99)
+      abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c
new file mode 100644
index 0000000..a11c647
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c
@@ -0,0 +1,83 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  const int S = 32;
+  mystruct *m = (mystruct *) calloc (S, sizeof (*m));
+  int i, j;
+
+  for (i = 0; i < S; i++)
+    {
+      m[i].a = (int *) malloc (N * sizeof (int));
+      m[i].b = (int *) malloc (N * sizeof (int));
+      m[i].c = (int *) malloc (N * sizeof (int));
+    }
+
+  for (j = 0; j < S; j++)
+    for (i = 0; i < N; i++)
+      {
+	m[j].a[i] = 0;
+	m[j].b[i] = 0;
+	m[j].c[i] = 0;
+      }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j, k;
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].a[0:N])
+        for (j = 0; j < N; j++)
+          m[k].a[j]++;
+
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10])
+	for (j = 0; j < N; j++)
+	  {
+	    m[k].b[j]++;
+	    if (j > 5 && j < N - 5)
+	      m[k].c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (j = 0; j < S; j++)
+    {
+      for (i = 0; i < N; i++)
+	{
+	  if (m[j].a[i] != 99)
+	    abort ();
+	  if (m[j].b[i] != 99)
+	    abort ();
+	  if (i > 5 && i < N-5)
+	    {
+	      if (m[j].c[i] != 99)
+		abort ();
+	    }
+	  else
+	    {
+	      if (m[j].c[i] != 0)
+		abort ();
+	    }
+	}
+
+      free (m[j].a);
+      free (m[j].b);
+      free (m[j].c);
+    }
+  free (m);
+
+  return 0;
+}