Fix gimple_copy for OpenMP atomic load/store [PR122281, PR105001]

	PR libgomp/122281
	PR middle-end/105001

gcc/ChangeLog:

	* gimple.cc (gimple_copy): Add missing unshare_expr for
	GIMPLE_OMP_ATOMIC_LOAD and GIMPLE_OMP_ATOMIC_STORE.

diff --git a/gcc/gimple.cc b/gcc/gimple.cc
index 102e21fe5e5b2b82fa7c31383597eef9f3037949..b968a45aaa06000b5b487c21d04ae840575a1362 100644 (file)
--- a/gcc/gimple.cc
+++ b/gcc/gimple.cc
@@ -2283,6 +2283,28 @@ gimple_copy (gimple *stmt)
        }
     }
 
+  switch (gimple_code (stmt))
+    {
+    case GIMPLE_OMP_ATOMIC_LOAD:
+      {
+       gomp_atomic_load *g = as_a <gomp_atomic_load *> (copy);
+       gimple_omp_atomic_load_set_lhs (g,
+         unshare_expr (gimple_omp_atomic_load_lhs (g)));
+       gimple_omp_atomic_load_set_rhs (g,
+         unshare_expr (gimple_omp_atomic_load_rhs (g)));
+       break;
+      }
+    case GIMPLE_OMP_ATOMIC_STORE:
+      {
+       gomp_atomic_store *g = as_a <gomp_atomic_store *> (copy);
+       gimple_omp_atomic_store_set_val (g,
+         unshare_expr (gimple_omp_atomic_store_val (g)));
+       break;
+      }
+    default:
+      break;
+    }
+
   /* Make copy of operands.  */
   for (i = 0; i < num_ops; i++)
     gimple_set_op (copy, i, unshare_expr (gimple_op (stmt, i)));

diff --git a/libgomp/testsuite/libgomp.c/pr122281.c b/libgomp/testsuite/libgomp.c/pr122281.c
new file mode 100644 (file)
index 0000000..a02a728
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr122281.c
@@ -0,0 +1,43 @@
+/* { dg-do run }  */
+/* { dg-additional-options "-O3" }  */
+
+/* PR libgomp/122281  */
+/* PR middle-end/105001  */
+
+/* If SIMT is supported, the inner 'omp simd' is duplicated into
+   one SIMT and one SIMD variant. SIMT is currently only supported
+   with nvidia GPUs.  (This only happens with -O1 or higher.)
+
+   The duplication failed for the SIMD case as a tree was shared and
+   the initialization only happened in the SIMT branch, i.e. when
+   compiling for a SIMT-device, all non-SIMD (offload or host devices)
+   accesses failed (segfault) for the atomic update.  */
+
+#include <omp.h>
+
+int __attribute__((noinline, noclone))
+f(int *A, int n, int dev) {
+ int cnt = 0;
+ #pragma omp target map(cnt) device(dev)
+ {
+   #pragma omp parallel for simd
+   for (int i = 0; i < n; i++)
+   if (A[i] != 0)
+     {
+       #pragma omp atomic
+       cnt++;
+     }
+ }
+ return cnt;
+}
+
+int main() {
+  int n = 10;
+  int A[10] = {11,22,33,44,55,66,77,88,99,110};
+
+  /* Run over all devices, including the host; the host should be SIMD,
+     some non-host devices might be SIMT.  */
+  for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+    if (f (A, n, dev) != 10)
+      __builtin_abort();
+}