public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Handle references in OpenACC "private" clauses
@ 2021-05-13 16:12 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:12 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:20be384324445e7f620c081748a9cd77981705e4
commit 20be384324445e7f620c081748a9cd77981705e4
Author: Julian Brown <julian@codesourcery.com>
Date: Fri Sep 20 13:53:10 2019 -0700
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.
Diff:
---
gcc/ChangeLog.omp | 5 ++
gcc/gimplify.c | 15 +++++
libgomp/ChangeLog.omp | 6 ++
.../testsuite/libgomp.oacc-c++/privatized-ref-2.C | 64 +++++++++++++++++++
.../testsuite/libgomp.oacc-c++/privatized-ref-3.C | 64 +++++++++++++++++++
.../libgomp.oacc-fortran/privatized-ref-1.f95 | 71 ++++++++++++++++++++++
6 files changed, 225 insertions(+)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 72797d99687..63dc5b23904 100644
--- 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>
* gimplify.c (gomp_oacc_needs_data_present): Remove unused variables.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index e463511344c..f2c55cca5fb 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -11679,6 +11679,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 ff5b8dc785d..e0a6bff168f 100644
--- 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-17 Julian Brown <julian@codesourcery.com>
* oacc-host.c (host_openacc_async_queue_callback): Invoke callback
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
index 00000000000..3884f163132
--- /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
index 00000000000..c1a10cba31b
--- /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
index 00000000000..f16f69c1d1b
--- /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
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:12 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:12 [gcc/devel/omp/gcc-11] Handle references in OpenACC "private" clauses Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).