[og9] Handle references in OpenACC "private" clauses

	gcc/
	* gimplify.c (localize_reductions): Rewrite references for
	OMP_CLAUSE_PRIVATE also.

	libgomp/
	* testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: New test.
	* testsuite/libgomp.oacc-c++/privatized-ref-2.C: New test.
	* testsuite/libgomp.oacc-c++/privatized-ref-3.C: New test.

(cherry picked from openacc-gcc-9-branch commit
28b843e3c547bda9857d2c60ffcc44f343e8ac4f)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index fe5849591530d61486abfc46519565abf0c9c888..523b6eb1d747a2cbcc1d4077f073ca6c95cfec42 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2019-09-20  Julian Brown  <julian@codesourcery.com>
+
+       * gimplify.c (localize_reductions): Rewrite references for
+       OMP_CLAUSE_PRIVATE also.
+
 2019-09-17  Tobias Burnus  <tobias@codesourcery.com>
 
        * config/gcn/gcn.c (gcn_expand_scalar_to_vector_address,

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index d16611d3617e6f4ccdb1f26a9d42a497f0154bcd..d95ad5d4baa3d2fdcd6a24e481357610ff99d0c9 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10879,6 +10879,21 @@ localize_reductions (tree clauses, tree body)
 
        OMP_CLAUSE_REDUCTION_PRIVATE_DECL (c) = new_var;
       }
+    else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE)
+      {
+       var = OMP_CLAUSE_DECL (c);
+
+       if (!lang_hooks.decls.omp_privatize_by_reference (var))
+         continue;
+
+       type = TREE_TYPE (TREE_TYPE (var));
+       new_var = create_tmp_var (type, IDENTIFIER_POINTER (DECL_NAME (var)));
+
+       pr.ref_var = var;
+       pr.local_var = new_var;
+
+       walk_tree (&body, localize_reductions_r, &pr, NULL);
+      }
 }
 
 

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 7813760e64253472b166ad04dda5da740d7c36e2..d9d1c353e31ec1170d4401427af25cae6262aac5 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,9 @@
+2019-09-20  Julian Brown  <julian@codesourcery.com>
+
+       * testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: New test.
+       * testsuite/libgomp.oacc-c++/privatized-ref-2.C: New test.
+       * testsuite/libgomp.oacc-c++/privatized-ref-3.C: New test.
+
 2019-09-19  Julian Brown  <julian@codesourcery.com>
 
        * plugin/plugin-nvptx.c (GOMP_OFFLOAD_openacc_async_host2dev):

diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
new file mode 100644 (file)
index 0000000..3884f16
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+void workers (void)
+{
+  double res[65536];
+  int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
+  {
+    int i, j;
+#pragma acc loop gang
+    for (i = 0; i < 256; i++)
+      {
+#pragma acc loop worker
+       for (j = 0; j < 256; j++)
+         {
+           int tmpvar;
+           int &tmpref = tmpvar;
+           tmpref = (i * 256 + j) * 99;
+           res[i * 256 + j] = tmpref;
+         }
+      }
+  }
+
+  for (i = 0; i < 65536; i++)
+    if (res[i] != i * 99)
+      abort ();
+}
+
+void vectors (void)
+{
+  double res[65536];
+  int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
+  {
+    int i, j;
+#pragma acc loop gang worker
+    for (i = 0; i < 256; i++)
+      {
+#pragma acc loop vector
+       for (j = 0; j < 256; j++)
+         {
+           int tmpvar;
+           int &tmpref = tmpvar;
+           tmpref = (i * 256 + j) * 101;
+           res[i * 256 + j] = tmpref;
+         }
+      }
+  }
+
+  for (i = 0; i < 65536; i++)
+    if (res[i] != i * 101)
+      abort ();
+}
+
+int main (int argc, char *argv[])
+{
+  workers ();
+  vectors ();
+  return 0;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C
new file mode 100644 (file)
index 0000000..c1a10cb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+void workers (void)
+{
+  double res[65536];
+  int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
+  {
+    int i, j;
+    int tmpvar;
+    int &tmpref = tmpvar;
+#pragma acc loop gang
+    for (i = 0; i < 256; i++)
+      {
+#pragma acc loop worker private(tmpref)
+       for (j = 0; j < 256; j++)
+         {
+           tmpref = (i * 256 + j) * 99;
+           res[i * 256 + j] = tmpref;
+         }
+      }
+  }
+
+  for (i = 0; i < 65536; i++)
+    if (res[i] != i * 99)
+      abort ();
+}
+
+void vectors (void)
+{
+  double res[65536];
+  int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
+  {
+    int i, j;
+    int tmpvar;
+    int &tmpref = tmpvar;
+#pragma acc loop gang worker
+    for (i = 0; i < 256; i++)
+      {
+#pragma acc loop vector private(tmpref)
+       for (j = 0; j < 256; j++)
+         {
+           tmpref = (i * 256 + j) * 101;
+           res[i * 256 + j] = tmpref;
+         }
+      }
+  }
+
+  for (i = 0; i < 65536; i++)
+    if (res[i] != i * 101)
+      abort ();
+}
+
+int main (int argc, char *argv[])
+{
+  workers ();
+  vectors ();
+  return 0;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95
new file mode 100644 (file)
index 0000000..f16f69c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95
@@ -0,0 +1,71 @@
+! { dg-do run }
+
+program main
+  implicit none
+  integer :: myint
+  integer :: i
+  real :: res(65536), tmp
+
+  res(:) = 0.0
+
+  myint = 5
+  call workers(myint, res)
+
+  do i=1,65536
+    tmp = i * 99
+    if (res(i) .ne. tmp) stop 1
+  end do
+
+  res(:) = 0.0
+
+  myint = 7
+  call vectors(myint, res)
+
+  do i=1,65536
+    tmp = i * 101
+    if (res(i) .ne. tmp) stop 2
+  end do
+
+contains
+
+  subroutine workers(t1, res)
+    implicit none
+    integer :: t1
+    integer :: i, j
+    real, intent(out) :: res(:)
+
+    !$acc parallel copyout(res) num_gangs(64) num_workers(64)
+
+    !$acc loop gang
+    do i=0,255
+      !$acc loop worker private(t1)
+      do j=1,256
+        t1 = (i * 256 + j) * 99
+        res(i * 256 + j) = t1
+      end do
+    end do
+
+    !$acc end parallel
+  end subroutine workers
+
+  subroutine vectors(t1, res)
+    implicit none
+    integer :: t1
+    integer :: i, j
+    real, intent(out) :: res(:)
+
+    !$acc parallel copyout(res) num_gangs(64) num_workers(64)
+
+    !$acc loop gang worker
+    do i=0,255
+      !$acc loop vector private(t1)
+      do j=1,256
+        t1 = (i * 256 + j) * 101
+        res(i * 256 + j) = t1
+      end do
+    end do
+
+    !$acc end parallel
+  end subroutine vectors
+
+end program main