Merge commit 'a0c06cc27d2146b7d86758ffa236516c6143d62c' into devel/omp/gcc-9

This merges all patches after releases/gcc-9.1.0 until and including
to the commit belonging to the tag 'releases/gcc-9.2.0'. That's
r9-7613-ga0c06cc27d2146b7d86758ffa236516c6143d62c on the releases/gcc-9
branch.

10 files changed, 328 insertions, 32 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ff585e9..958572a 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,55 @@
+2019-08-12  Release Manager
+
+	* GCC 9.2.0 released.
+
+2019-08-01  Jakub Jelinek  <jakub@redhat.com>
+
+	Backported from mainline
+	2019-07-31  Jakub Jelinek  <jakub@redhat.com>
+
+	PR middle-end/91301
+	* testsuite/libgomp.c++/for-27.C: New test.
+
+2019-06-11  Jakub Jelinek  <jakub@redhat.com>
+
+	PR target/90811
+	* testsuite/libgomp.c/pr90811.c: New test.
+
+2019-06-05  Jakub Jelinek  <jakub@redhat.com>
+
+	Backported from mainline
+	2019-05-27  Jakub Jelinek  <jakub@redhat.com>
+
+	PR libgomp/90641
+	* work.c (gomp_init_work_share): Instead of aligning final ordered
+	value to multiples of long long alignment, align to that the
+	first part (ordered team ids) and if inline_ordered_team_ids
+	is not on a long long alignment boundary within the structure,
+	use __alignof__ (long long) - 1 pad size always.
+	* loop.c (GOMP_loop_start): Fix *mem computation if
+	inline_ordered_team_ids is not aligned on long long alignment boundary
+	within the structure.
+	* loop-ull.c (GOMP_loop_ull_start): Likewise.
+	* sections.c (GOMP_sections2_start): Likewise.
+
+	2019-05-24  Jakub Jelinek  <jakub@redhat.com>
+
+	PR libgomp/90585
+	* plugin/plugin-hsa.c (print_kernel_dispatch, run_kernel): Use PRIu64
+	macro instead of "lu".
+	(release_kernel_dispatch): Likewise.  Cast shadow->debug to uintptr_t
+	before casting to void *.
+
+2019-05-20  Jakub Jelinek  <jakub@redhat.com>
+
+	PR libgomp/90527
+	* alloc.c (_GNU_SOURCE): Define.
+
+2019-05-07  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR target/87835
+	* testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update.
+
 2019-05-03  Release Manager
 
 	* GCC 9.1.0 released.
diff --git a/libgomp/alloc.c b/libgomp/alloc.c
index 9145e4b..c28e32f 100644
--- a/libgomp/alloc.c
+++ b/libgomp/alloc.c
@@ -27,6 +27,7 @@
    places in the OpenMP API do not make any provision for failure, so in
    general we cannot allow memory allocation to fail.  */
 
+#define _GNU_SOURCE
 #include "libgomp.h"
 #include <stdlib.h>
 
diff --git a/libgomp/loop.c b/libgomp/loop.c
index 7248cb5..37861d1 100644
--- a/libgomp/loop.c
+++ b/libgomp/loop.c
@@ -267,14 +267,17 @@ GOMP_loop_start (long start, long end, long incr, long sched,
       if (mem)
 	{
 	  uintptr_t size = (uintptr_t) *mem;
+#define INLINE_ORDERED_TEAM_IDS_OFF \
+  ((offsetof (struct gomp_work_share, inline_ordered_team_ids)		\
+    + __alignof__ (long long) - 1) & ~(__alignof__ (long long) - 1))
 	  if (size > (sizeof (struct gomp_work_share)
-		      - offsetof (struct gomp_work_share,
-				  inline_ordered_team_ids)))
-	    thr->ts.work_share->ordered_team_ids
-	      = gomp_malloc_cleared (size);
+		      - INLINE_ORDERED_TEAM_IDS_OFF))
+	    *mem
+	      = (void *) (thr->ts.work_share->ordered_team_ids
+			  = gomp_malloc_cleared (size));
 	  else
-	    memset (thr->ts.work_share->ordered_team_ids, '\0', size);
-	  *mem = (void *) thr->ts.work_share->ordered_team_ids;
+	    *mem = memset (((char *) thr->ts.work_share)
+			   + INLINE_ORDERED_TEAM_IDS_OFF, '\0', size);
 	}
       gomp_work_share_init_done ();
     }
@@ -287,7 +290,18 @@ GOMP_loop_start (long start, long end, long incr, long sched,
 						  first_reductions);
 	}
       if (mem)
-	*mem = (void *) thr->ts.work_share->ordered_team_ids;
+	{
+	  if ((offsetof (struct gomp_work_share, inline_ordered_team_ids)
+	       & (__alignof__ (long long) - 1)) == 0)
+	    *mem = (void *) thr->ts.work_share->ordered_team_ids;
+	  else
+	    {
+	      uintptr_t p = (uintptr_t) thr->ts.work_share->ordered_team_ids;
+	      p += __alignof__ (long long) - 1;
+	      p &= ~(__alignof__ (long long) - 1);
+	      *mem = (void *) p;
+	    }
+	}
     }
 
   if (!istart)
diff --git a/libgomp/loop_ull.c b/libgomp/loop_ull.c
index 502bb5e..edf1c32 100644
--- a/libgomp/loop_ull.c
+++ b/libgomp/loop_ull.c
@@ -266,14 +266,17 @@ GOMP_loop_ull_start (bool up, gomp_ull start, gomp_ull end,
       if (mem)
 	{
 	  uintptr_t size = (uintptr_t) *mem;
+#define INLINE_ORDERED_TEAM_IDS_OFF \
+  ((offsetof (struct gomp_work_share, inline_ordered_team_ids)		\
+    + __alignof__ (long long) - 1) & ~(__alignof__ (long long) - 1))
 	  if (size > (sizeof (struct gomp_work_share)
-		      - offsetof (struct gomp_work_share,
-				  inline_ordered_team_ids)))
-	    thr->ts.work_share->ordered_team_ids
-	      = gomp_malloc_cleared (size);
+		      - INLINE_ORDERED_TEAM_IDS_OFF))
+	    *mem
+	      = (void *) (thr->ts.work_share->ordered_team_ids
+			  = gomp_malloc_cleared (size));
 	  else
-	    memset (thr->ts.work_share->ordered_team_ids, '\0', size);
-	  *mem = (void *) thr->ts.work_share->ordered_team_ids;
+	    *mem = memset (((char *) thr->ts.work_share)
+			   + INLINE_ORDERED_TEAM_IDS_OFF, '\0', size);
 	}
       gomp_work_share_init_done ();
     }
@@ -286,7 +289,18 @@ GOMP_loop_ull_start (bool up, gomp_ull start, gomp_ull end,
 						  first_reductions);
 	}
       if (mem)
-	*mem = (void *) thr->ts.work_share->ordered_team_ids;
+	{
+	  if ((offsetof (struct gomp_work_share, inline_ordered_team_ids)
+	       & (__alignof__ (long long) - 1)) == 0)
+	    *mem = (void *) thr->ts.work_share->ordered_team_ids;
+	  else
+	    {
+	      uintptr_t p = (uintptr_t) thr->ts.work_share->ordered_team_ids;
+	      p += __alignof__ (long long) - 1;
+	      p &= ~(__alignof__ (long long) - 1);
+	      *mem = (void *) p;
+	    }
+	}
     }
 
   return ialias_call (GOMP_loop_ull_runtime_next) (istart, iend);
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index 3c990c5..8b4dbfb 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -1180,8 +1180,9 @@ create_single_kernel_dispatch (struct kernel_info *kernel,
 static void
 release_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *shadow)
 {
-  HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n", shadow,
-	     shadow->debug, (void *) shadow->debug);
+  HSA_DEBUG ("Released kernel dispatch: %p has value: %" PRIu64 " (%p)\n",
+	     shadow, shadow->debug,
+	     (void *) (uintptr_t) shadow->debug);
 
   hsa_fns.hsa_memory_free_fn (shadow->kernarg_address);
 
@@ -1302,9 +1303,9 @@ print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned inden
   indent_stream (stderr, indent);
   fprintf (stderr, "kernarg_address: %p\n", dispatch->kernarg_address);
   indent_stream (stderr, indent);
-  fprintf (stderr, "object: %lu\n", dispatch->object);
+  fprintf (stderr, "object: %" PRIu64 "\n", dispatch->object);
   indent_stream (stderr, indent);
-  fprintf (stderr, "signal: %lu\n", dispatch->signal);
+  fprintf (stderr, "signal: %" PRIu64 "\n", dispatch->signal);
   indent_stream (stderr, indent);
   fprintf (stderr, "private_segment_size: %u\n",
 	   dispatch->private_segment_size);
@@ -1312,7 +1313,7 @@ print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned inden
   fprintf (stderr, "group_segment_size: %u\n",
 	   dispatch->group_segment_size);
   indent_stream (stderr, indent);
-  fprintf (stderr, "children dispatches: %lu\n",
+  fprintf (stderr, "children dispatches: %" PRIu64 "\n",
 	   dispatch->kernel_dispatch_count);
   indent_stream (stderr, indent);
   fprintf (stderr, "omp_num_threads: %u\n",
@@ -1620,7 +1621,7 @@ run_kernel (struct kernel_info *kernel, void *vars,
 	hsa_signal_t child_s;
 	child_s.handle = shadow->children_dispatches[i]->signal;
 
-	HSA_DEBUG ("Waiting for children completion signal: %lu\n",
+	HSA_DEBUG ("Waiting for children completion signal: %" PRIu64 "\n",
 		   shadow->children_dispatches[i]->signal);
 	hsa_fns.hsa_signal_load_acquire_fn (child_s);
       }
diff --git a/libgomp/sections.c b/libgomp/sections.c
index bf7baae..8879436 100644
--- a/libgomp/sections.c
+++ b/libgomp/sections.c
@@ -118,14 +118,17 @@ GOMP_sections2_start (unsigned count, uintptr_t *reductions, void **mem)
       if (mem)
 	{
 	  uintptr_t size = (uintptr_t) *mem;
+#define INLINE_ORDERED_TEAM_IDS_OFF \
+  ((offsetof (struct gomp_work_share, inline_ordered_team_ids)		\
+    + __alignof__ (long long) - 1) & ~(__alignof__ (long long) - 1))
 	  if (size > (sizeof (struct gomp_work_share)
-		      - offsetof (struct gomp_work_share,
-				  inline_ordered_team_ids)))
-	    thr->ts.work_share->ordered_team_ids
-	      = gomp_malloc_cleared (size);
+		      - INLINE_ORDERED_TEAM_IDS_OFF))
+	    *mem
+	      = (void *) (thr->ts.work_share->ordered_team_ids
+			  = gomp_malloc_cleared (size));
 	  else
-	    memset (thr->ts.work_share->ordered_team_ids, '\0', size);
-	  *mem = (void *) thr->ts.work_share->ordered_team_ids;
+	    *mem = memset (((char *) thr->ts.work_share)
+			   + INLINE_ORDERED_TEAM_IDS_OFF, '\0', size);
 	}
       gomp_work_share_init_done ();
     }
@@ -138,7 +141,18 @@ GOMP_sections2_start (unsigned count, uintptr_t *reductions, void **mem)
 						  first_reductions);
 	}
       if (mem)
-	*mem = (void *) thr->ts.work_share->ordered_team_ids;
+	{
+	  if ((offsetof (struct gomp_work_share, inline_ordered_team_ids)
+	       & (__alignof__ (long long) - 1)) == 0)
+	    *mem = (void *) thr->ts.work_share->ordered_team_ids;
+	  else
+	    {
+	      uintptr_t p = (uintptr_t) thr->ts.work_share->ordered_team_ids;
+	      p += __alignof__ (long long) - 1;
+	      p &= ~(__alignof__ (long long) - 1);
+	      *mem = (void *) p;
+	    }
+	}
     }
 
 #ifdef HAVE_SYNC_BUILTINS
diff --git a/libgomp/testsuite/libgomp.c++/for-27.C b/libgomp/testsuite/libgomp.c++/for-27.C
new file mode 100644
index 0000000..7dca430
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/for-27.C
@@ -0,0 +1,169 @@
+// { dg-do run }
+
+typedef __PTRDIFF_TYPE__ ptrdiff_t;
+extern "C" void abort ();
+
+int a[2000];
+
+template <typename T>
+class I
+{
+public:
+  typedef ptrdiff_t difference_type;
+  I ();
+  ~I ();
+  I (T *);
+  I (const I &);
+  T &operator * ();
+  T *operator -> ();
+  T &operator [] (const difference_type &) const;
+  I &operator = (const I &);
+  I &operator ++ ();
+  I operator ++ (int);
+  I &operator -- ();
+  I operator -- (int);
+  I &operator += (const difference_type &);
+  I &operator -= (const difference_type &);
+  I operator + (const difference_type &) const;
+  I operator - (const difference_type &) const;
+  template <typename S> friend bool operator == (I<S> &, I<S> &);
+  template <typename S> friend bool operator == (const I<S> &, const I<S> &);
+  template <typename S> friend bool operator < (I<S> &, I<S> &);
+  template <typename S> friend bool operator < (const I<S> &, const I<S> &);
+  template <typename S> friend bool operator <= (I<S> &, I<S> &);
+  template <typename S> friend bool operator <= (const I<S> &, const I<S> &);
+  template <typename S> friend bool operator > (I<S> &, I<S> &);
+  template <typename S> friend bool operator > (const I<S> &, const I<S> &);
+  template <typename S> friend bool operator >= (I<S> &, I<S> &);
+  template <typename S> friend bool operator >= (const I<S> &, const I<S> &);
+  template <typename S> friend typename I<S>::difference_type operator - (I<S> &, I<S> &);
+  template <typename S> friend typename I<S>::difference_type operator - (const I<S> &, const I<S> &);
+  template <typename S> friend I<S> operator + (typename I<S>::difference_type , const I<S> &);
+private:
+  T *p;
+};
+template <typename T> I<T>::I () : p (0) {}
+template <typename T> I<T>::~I () {}
+template <typename T> I<T>::I (T *x) : p (x) {}
+template <typename T> I<T>::I (const I &x) : p (x.p) {}
+template <typename T> T &I<T>::operator * () { return *p; }
+template <typename T> T *I<T>::operator -> () { return p; }
+template <typename T> T &I<T>::operator [] (const difference_type &x) const { return p[x]; }
+template <typename T> I<T> &I<T>::operator = (const I &x) { p = x.p; return *this; }
+template <typename T> I<T> &I<T>::operator ++ () { ++p; return *this; }
+template <typename T> I<T> I<T>::operator ++ (int) { return I (p++); }
+template <typename T> I<T> &I<T>::operator -- () { --p; return *this; }
+template <typename T> I<T> I<T>::operator -- (int) { return I (p--); }
+template <typename T> I<T> &I<T>::operator += (const difference_type &x) { p += x; return *this; }
+template <typename T> I<T> &I<T>::operator -= (const difference_type &x) { p -= x; return *this; }
+template <typename T> I<T> I<T>::operator + (const difference_type &x) const { return I (p + x); }
+template <typename T> I<T> I<T>::operator - (const difference_type &x) const { return I (p - x); }
+template <typename T> bool operator == (I<T> &x, I<T> &y) { return x.p == y.p; }
+template <typename T> bool operator == (const I<T> &x, const I<T> &y) { return x.p == y.p; }
+template <typename T> bool operator != (I<T> &x, I<T> &y) { return !(x == y); }
+template <typename T> bool operator != (const I<T> &x, const I<T> &y) { return !(x == y); }
+template <typename T> bool operator < (I<T> &x, I<T> &y) { return x.p < y.p; }
+template <typename T> bool operator < (const I<T> &x, const I<T> &y) { return x.p < y.p; }
+template <typename T> bool operator <= (I<T> &x, I<T> &y) { return x.p <= y.p; }
+template <typename T> bool operator <= (const I<T> &x, const I<T> &y) { return x.p <= y.p; }
+template <typename T> bool operator > (I<T> &x, I<T> &y) { return x.p > y.p; }
+template <typename T> bool operator > (const I<T> &x, const I<T> &y) { return x.p > y.p; }
+template <typename T> bool operator >= (I<T> &x, I<T> &y) { return x.p >= y.p; }
+template <typename T> bool operator >= (const I<T> &x, const I<T> &y) { return x.p >= y.p; }
+template <typename T> typename I<T>::difference_type operator - (I<T> &x, I<T> &y) { return x.p - y.p; }
+template <typename T> typename I<T>::difference_type operator - (const I<T> &x, const I<T> &y) { return x.p - y.p; }
+template <typename T> I<T> operator + (typename I<T>::difference_type x, const I<T> &y) { return I<T> (x + y.p); }
+
+template <typename T>
+class J
+{
+public:
+  J(const I<T> &x, const I<T> &y) : b (x), e (y) {}
+  const I<T> &begin ();
+  const I<T> &end ();
+private:
+  I<T> b, e;
+};
+
+template <typename T> const I<T> &J<T>::begin () { return b; }
+template <typename T> const I<T> &J<T>::end () { return e; }
+
+int results[2000];
+
+template <typename T> 
+void
+baz (I<T> &i)
+{
+  if (*i < 0 || *i >= 2000)
+    abort ();
+  results[*i]++;
+}
+
+static inline void
+baz (int i)
+{
+  results[i]++;
+}
+
+void
+f1 ()
+{
+#pragma omp simd
+  for (auto i : a)
+    baz (i);
+}
+
+void
+f2 (const I<int> &x, const I<int> &y)
+{
+  I<int> i;
+#pragma omp distribute parallel for
+  for (i = x; i <= y; i += 6)
+    baz (*i);
+}
+
+void
+f3 (const I<int> &x, const I<int> &y)
+{
+  I<int> i;
+#pragma omp distribute parallel for private (i)
+  for (i = x; i < y - 1; i = 1 - 6 + 7 + i)
+    baz (*i);
+}
+
+void
+f4 (const I<int> &x, const I<int> &y)
+{
+  I<int> i;
+#pragma omp teams distribute parallel for lastprivate (i)
+  for (i = x + 2000 - 64; i > y + 10; --i)
+    baz (*i);
+}
+
+#define check(expr) \
+  for (int i = 0; i < 2000; i++)			\
+    if (expr)						\
+      {							\
+	if (results[i] != 1)				\
+	  abort ();					\
+	results[i] = 0;					\
+      }							\
+    else if (results[i])				\
+      abort ()
+
+int
+main ()
+{
+  for (int i = 0; i < 2000; i++)
+    a[i] = i;
+  f1 ();
+  check (1);
+  #pragma omp teams
+  f2 (&a[10], &a[1990]);
+  check (i >= 10 && i <= 1990 && (i - 10) % 6 == 0);
+  #pragma omp teams
+  f3 (&a[0], &a[1999]);
+  check (i < 1998 && (i & 1) == 0);
+  f4 (&a[0], &a[30]);
+  check (i > 40 && i <= 2000 - 64);
+}
diff --git a/libgomp/testsuite/libgomp.c/pr90811.c b/libgomp/testsuite/libgomp.c/pr90811.c
new file mode 100644
index 0000000..25b7d78
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr90811.c
@@ -0,0 +1,29 @@
+/* PR target/90811 */
+
+int
+main ()
+{
+  long long a[100], b[100];
+  int i;
+  for (i = 0; i < 100; i++)
+    {
+      a[i] = i;
+      b[i] = i % 10;
+    }
+  #pragma omp target teams distribute parallel for simd map(tofrom: a[:100], b[:100])
+  for (i = 0; i < 100; i++)
+    {
+      long long c = 0;
+      const long long d[] = { 1, 3, 5, 7, 9 };
+      for (int j = 4; j >= 0; j--)
+         c = d[j] + b[i] * c;
+      a[i] += c;
+    }
+  for (i = 0; i < 100; i++)
+    {
+      const long long r[] = { 1, 26, 229, 976, 2849, 6646, 13381, 24284, 40801, 64594 };
+      if (a[i] != r[i % 10] + (i / 10 * 10))
+	__builtin_abort ();
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c
index 503ce17..88c2c77 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c
@@ -16,7 +16,6 @@ main (void)
   CUstream stream1;
   int N = n;
   int a[n];
-  int b[n];
   int c[n];
 
   acc_init (acc_device_nvidia);
@@ -36,7 +35,7 @@ main (void)
       c[i] = 0;
     }
 
-#pragma acc data copy (a, b, c) copyin (N)
+#pragma acc data copy (a, c) copyin (N)
   {
 #pragma acc parallel async (1)
     ;
diff --git a/libgomp/work.c b/libgomp/work.c
index 22adb8b..a589b8b 100644
--- a/libgomp/work.c
+++ b/libgomp/work.c
@@ -110,9 +110,12 @@ gomp_init_work_share (struct gomp_work_share *ws, size_t ordered,
 
       if (__builtin_expect (ordered != 1, 0))
 	{
-	  ordered += nthreads * sizeof (*ws->ordered_team_ids) - 1;
-	  ordered = ordered + __alignof__ (long long) - 1;
-	  ordered &= ~(__alignof__ (long long) - 1);
+	  size_t o = nthreads * sizeof (*ws->ordered_team_ids);
+	  o += __alignof__ (long long) - 1;
+	  if ((offsetof (struct gomp_work_share, inline_ordered_team_ids)
+	       & (__alignof__ (long long) - 1)) == 0)
+	    o &= ~(__alignof__ (long long) - 1);
+	  ordered += o - 1;
 	}
       else
 	ordered = nthreads * sizeof (*ws->ordered_team_ids);