Merge branches/gcc-5-branch rev 227732.linaro-local/gcc-5-integration-branch-new

Change-Id: I2f59904b28323b1c72a8cf1bd62c9e460d95bcea

13 files changed, 232 insertions, 10 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 6c23ff19f1d..d59d80b4773 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,55 @@
+2015-08-24  Joost VandeVondele  <vondele@gnu.gcc.org>
+
+	PR libgomp/66761
+	PR libgomp/67303
+	* iter.c (gomp_iter_dynamic_next): Employ an atomic load.
+	(gomp_iter_guided_next): Idem.
+	* iter_ull.c (gomp_iter_ull_dynamic_next): Idem.
+	(gomp_iter_ull_guided_next): Idem.
+	* config/linux/wait.h (do_spin): Idem.
+
+2015-07-16  Release Manager
+
+	* GCC 5.2.0 released.
+
+2015-07-15  Thomas Schwinge  <thomas@codesourcery.com>
+
+	Backport trunk r225560:
+
+	2015-07-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR libgomp/65099
+	* plugin/plugin-nvptx.c (nvptx_get_num_devices): Return 0 if not
+	in a 64-bit configuration.
+	* testsuite/libgomp.oacc-c++/c++.exp: Don't attempt nvidia
+	offloading testing if no such device is available.
+	* testsuite/libgomp.oacc-c/c.exp: Likewise.
+	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+
+	Backport trunk r223801:
+
+	2015-05-28  Julian Brown  <julian@codesourcery.com>
+
+	PR libgomp/65742
+
+	* oacc-init.c (plugin/plugin-host.h): Include.
+	(acc_on_device): Check whether we're in an offloaded region for
+	host_nonshm
+	plugin. Don't use __builtin_acc_on_device.
+	* plugin/plugin-host.c (GOMP_OFFLOAD_openacc_parallel): Set
+	nonshm_exec flag in thread-local data.
+	(GOMP_OFFLOAD_openacc_create_thread_data): Allocate thread-local
+	data for host_nonshm plugin.
+	(GOMP_OFFLOAD_openacc_destroy_thread_data): Free thread-local data
+	for host_nonshm plugin.
+	* plugin/plugin-host.h: New.
+
+2015-06-30  Jakub Jelinek  <jakub@redhat.com>
+
+	PR middle-end/66702
+	* testsuite/libgomp.c++/pr66702-1.C: New test.
+	* testsuite/libgomp.c++/pr66702-2.C: New test.
+
 2015-05-22  Jakub Jelinek  <jakub@redhat.com>
 
 	Backported from mainline
diff --git a/libgomp/config/linux/wait.h b/libgomp/config/linux/wait.h
index 96d2fbe5f8c..46a46155353 100644
--- a/libgomp/config/linux/wait.h
+++ b/libgomp/config/linux/wait.h
@@ -49,7 +49,9 @@ static inline int do_spin (int *addr, int val)
 {
   unsigned long long i, count = gomp_spin_count_var;
 
-  if (__builtin_expect (gomp_managed_threads > gomp_available_cpus, 0))
+  if (__builtin_expect (__atomic_load_n (&gomp_managed_threads,
+                                         MEMMODEL_RELAXED)
+                        > gomp_available_cpus, 0))
     count = gomp_throttled_spin_count_var;
   for (i = 0; i < count; i++)
     if (__builtin_expect (__atomic_load_n (addr, MEMMODEL_RELAXED) != val, 0))
diff --git a/libgomp/iter.c b/libgomp/iter.c
index 0ceb41d909e..1def8bde066 100644
--- a/libgomp/iter.c
+++ b/libgomp/iter.c
@@ -218,7 +218,7 @@ gomp_iter_dynamic_next (long *pstart, long *pend)
 	}
     }
 
-  start = ws->next;
+  start = __atomic_load_n (&ws->next, MEMMODEL_RELAXED);
   while (1)
     {
       long left = end - start;
@@ -301,7 +301,7 @@ gomp_iter_guided_next (long *pstart, long *pend)
   long start, end, nend, incr;
   unsigned long chunk_size;
 
-  start = ws->next;
+  start = __atomic_load_n (&ws->next, MEMMODEL_RELAXED);
   end = ws->end;
   incr = ws->incr;
   chunk_size = ws->chunk_size;
diff --git a/libgomp/iter_ull.c b/libgomp/iter_ull.c
index b1cad84d4c8..1c2d118b631 100644
--- a/libgomp/iter_ull.c
+++ b/libgomp/iter_ull.c
@@ -219,7 +219,7 @@ gomp_iter_ull_dynamic_next (gomp_ull *pstart, gomp_ull *pend)
 	}
     }
 
-  start = ws->next_ull;
+  start = __atomic_load_n (&ws->next_ull, MEMMODEL_RELAXED);
   while (1)
     {
       gomp_ull left = end - start;
@@ -305,7 +305,7 @@ gomp_iter_ull_guided_next (gomp_ull *pstart, gomp_ull *pend)
   gomp_ull start, end, nend, incr;
   gomp_ull chunk_size;
 
-  start = ws->next_ull;
+  start = __atomic_load_n (&ws->next_ull, MEMMODEL_RELAXED);
   end = ws->end_ull;
   incr = ws->incr_ull;
   chunk_size = ws->chunk_size_ull;
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index dc40fb6ffe1..a7c2e0d8208 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -29,6 +29,7 @@
 #include "libgomp.h"
 #include "oacc-int.h"
 #include "openacc.h"
+#include "plugin/plugin-host.h"
 #include <assert.h>
 #include <stdlib.h>
 #include <strings.h>
@@ -548,11 +549,18 @@ ialias (acc_set_device_num)
 int
 acc_on_device (acc_device_t dev)
 {
-  if (acc_get_device_type () == acc_device_host_nonshm)
+  struct goacc_thread *thr = goacc_thread ();
+
+  /* We only want to appear to be the "host_nonshm" plugin from "offloaded"
+     code -- i.e. within a parallel region.  Test a flag set by the
+     openacc_parallel hook of the host_nonshm plugin to determine that.  */
+  if (acc_get_device_type () == acc_device_host_nonshm
+      && thr && thr->target_tls
+      && ((struct nonshm_thread *)thr->target_tls)->nonshm_exec)
     return dev == acc_device_host_nonshm || dev == acc_device_not_host;
 
-  /* Just rely on the compiler builtin.  */
-  return __builtin_acc_on_device (dev);
+  /* For OpenACC, libgomp is only built for the host, so this is sufficient.  */
+  return dev == acc_device_host || dev == acc_device_none;
 }
 
 ialias (acc_on_device)
diff --git a/libgomp/plugin/plugin-host.c b/libgomp/plugin/plugin-host.c
index 1faf5bc194e..3cb4dab3778 100644
--- a/libgomp/plugin/plugin-host.c
+++ b/libgomp/plugin/plugin-host.c
@@ -44,6 +44,7 @@
 #include <stdlib.h>
 #include <string.h>
 #include <stdio.h>
+#include <stdbool.h>
 
 #ifdef HOST_NONSHM_PLUGIN
 #define STATIC
@@ -55,6 +56,10 @@
 #define SELF "host: "
 #endif
 
+#ifdef HOST_NONSHM_PLUGIN
+#include "plugin-host.h"
+#endif
+
 STATIC const char *
 GOMP_OFFLOAD_get_name (void)
 {
@@ -174,7 +179,10 @@ GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *),
 			       void *targ_mem_desc __attribute__ ((unused)))
 {
 #ifdef HOST_NONSHM_PLUGIN
+  struct nonshm_thread *thd = GOMP_PLUGIN_acc_thread ();
+  thd->nonshm_exec = true;
   fn (devaddrs);
+  thd->nonshm_exec = false;
 #else
   fn (hostaddrs);
 #endif
@@ -232,11 +240,20 @@ STATIC void *
 GOMP_OFFLOAD_openacc_create_thread_data (int ord
 					 __attribute__ ((unused)))
 {
+#ifdef HOST_NONSHM_PLUGIN
+  struct nonshm_thread *thd
+    = GOMP_PLUGIN_malloc (sizeof (struct nonshm_thread));
+  thd->nonshm_exec = false;
+  return thd;
+#else
   return NULL;
+#endif
 }
 
 STATIC void
-GOMP_OFFLOAD_openacc_destroy_thread_data (void *tls_data
-					  __attribute__ ((unused)))
+GOMP_OFFLOAD_openacc_destroy_thread_data (void *tls_data)
 {
+#ifdef HOST_NONSHM_PLUGIN
+  free (tls_data);
+#endif
 }
diff --git a/libgomp/plugin/plugin-host.h b/libgomp/plugin/plugin-host.h
new file mode 100644
index 00000000000..96955d19414
--- /dev/null
+++ b/libgomp/plugin/plugin-host.h
@@ -0,0 +1,37 @@
+/* OpenACC Runtime Library: acc_device_host, acc_device_host_nonshm.
+
+   Copyright (C) 2015 Free Software Foundation, Inc.
+
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#ifndef PLUGIN_HOST_H
+#define PLUGIN_HOST_H
+
+struct nonshm_thread
+{
+  bool nonshm_exec;
+};
+
+#endif
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 583ec87aeee..f9754b98f5a 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -777,6 +777,11 @@ nvptx_get_num_devices (void)
   int n;
   CUresult r;
 
+  /* PR libgomp/65099: Currently, we only support offloading in 64-bit
+     configurations.  */
+  if (sizeof (void *) != 8)
+    return 0;
+
   /* This function will be called before the plugin has been initialized in
      order to enumerate available devices, but CUDA API routines can't be used
      until cuInit has been called.  Just call it now (but don't yet do any
diff --git a/libgomp/testsuite/libgomp.c++/pr66702-1.C b/libgomp/testsuite/libgomp.c++/pr66702-1.C
new file mode 100644
index 00000000000..15772561b1a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66702-1.C
@@ -0,0 +1,49 @@
+// PR middle-end/66702
+// { dg-options "-O2" }
+// { dg-additional-options "-msse2" { target sse2_runtime } }
+// { dg-additional-options "-mavx" { target avx_runtime } }
+
+void
+bar (int &a, int &b, int *&c, int &d)
+{
+  volatile int x;
+  int *volatile y;
+  x = a; a = x;
+  x = b; b = x;
+  y = c; c = y;
+  x = d; d = x;
+}
+
+void (*volatile barp) (int &, int &, int *&, int &) = bar;
+
+#pragma omp declare simd uniform(b, c) linear(d:2) aligned(c:32) notinbranch
+int
+foo (int a, int b, int *c, int d)
+{
+  a++;
+  b++;
+  c += 8;
+  d += 2;
+  barp (a, b, c, d);
+  return a + b + *c + d;
+}
+
+volatile int e = 5;
+int c[64] __attribute__((aligned (32)));
+
+int
+main ()
+{
+  int d = 7, r = 0;
+  int b = e;
+  for (int i = 0; i < 64; i++)
+    c[i] = i;
+  #pragma omp simd reduction(+:r) linear(d:2)
+  for (int i = 0; i < 64; i++)
+    {
+      r += foo (i, b, c, d);
+      d += 2;
+    }
+  if (r != 7584)
+    __builtin_abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/pr66702-2.C b/libgomp/testsuite/libgomp.c++/pr66702-2.C
new file mode 100644
index 00000000000..7de3de010a7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr66702-2.C
@@ -0,0 +1,34 @@
+// PR middle-end/66702
+// { dg-options "-O2" }
+// { dg-additional-options "-msse2" { target sse2_runtime } }
+// { dg-additional-options "-mavx" { target avx_runtime } }
+
+struct S { int s1, s2; };
+struct T { T (); ~T (); int t; };
+
+T::T () : t(0) {}
+T::~T () {}
+
+#pragma omp declare simd uniform(b, c) notinbranch
+__attribute__((noinline)) int
+foo (int a, S b, T c)
+{
+  a++;
+  b.s1++;
+  b.s2++;
+  c.t++;
+  return a + b.s1 + b.s2 + c.t;
+}
+
+int
+main ()
+{
+  int r = 0;
+  S s = { 2, 3 };
+  T t;
+  #pragma omp simd reduction(+:r)
+  for (int i = 0; i < 64; i++)
+    r += foo (i, s, t);
+  if (r != 2592)
+    __builtin_abort ();
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
index f486f9b97ba..8941432b3e8 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -80,6 +80,12 @@ if { $lang_test_file_found } {
 		set acc_mem_shared 0
 	    }
 	    nvidia {
+		if { ![check_effective_target_openacc_nvidia_accel_present] } {
+		    # Don't bother; execution testing is going to FAIL.
+		    untested "$subdir $offload_target_openacc offloading"
+		    continue
+		}
+
 		# Copy ptx file (TEMPORARY)
 		remote_download host $srcdir/libgomp.oacc-c-c++-common/subr.ptx
 
diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp
index c0c70bbacad..326b988d99a 100644
--- a/libgomp/testsuite/libgomp.oacc-c/c.exp
+++ b/libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -48,6 +48,12 @@ foreach offload_target_openacc $offload_targets_s_openacc {
 	    set acc_mem_shared 0
 	}
 	nvidia {
+	    if { ![check_effective_target_openacc_nvidia_accel_present] } {
+		# Don't bother; execution testing is going to FAIL.
+		untested "$subdir $offload_target_openacc offloading"
+		continue
+	    }
+
 	    # Copy ptx file (TEMPORARY)
 	    remote_download host $srcdir/libgomp.oacc-c-c++-common/subr.ptx
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index a8f62e84d67..a8aaff0e2b1 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -77,6 +77,12 @@ if { $lang_test_file_found } {
 		set acc_mem_shared 0
 	    }
 	    nvidia {
+		if { ![check_effective_target_openacc_nvidia_accel_present] } {
+		    # Don't bother; execution testing is going to FAIL.
+		    untested "$subdir $offload_target_openacc offloading"
+		    continue
+		}
+
 		set acc_mem_shared 0
 	    }
 	    default {