From: Frederik Harwath <frederik@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: Julian Brown <julian@codesourcery.com>, <thomas@codesourcery.com>
Subject: [PATCH 35/40] Handle references in OpenACC "private" clauses
Date: Wed, 15 Dec 2021 16:54:42 +0100 [thread overview]
Message-ID: <20211215155447.19379-36-frederik@codesourcery.com> (raw)
In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com>
From: Julian Brown <julian@codesourcery.com>
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.
---
gcc/gimplify.c | 15 ++++
.../libgomp.oacc-c++/privatized-ref-2.C | 64 +++++++++++++++++
.../libgomp.oacc-c++/privatized-ref-3.C | 64 +++++++++++++++++
.../libgomp.oacc-fortran/privatized-ref-1.f95 | 71 +++++++++++++++++++
4 files changed, 214 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
create mode 100644 libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index daa69ccf6202..bf37388f947c 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -11976,6 +11976,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/testsuite/libgomp.oacc-c++/privatized-ref-2.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
new file mode 100644
index 000000000000..3884f163132c
--- /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 000000000000..c1a10cba31b3
--- /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 000000000000..fe1520a8078c
--- /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) ! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } }
+
+ !$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) ! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } }
+
+ !$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
--
2.33.0
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
next prev parent reply other threads:[~2021-12-15 15:57 UTC|newest]
Thread overview: 49+ messages / expand[flat|nested] mbox.gz Atom feed top
2021-12-15 15:54 [PATCH 00/40] OpenACC "kernels" Improvements Frederik Harwath
2021-12-15 15:54 ` [PATCH 01/40] Kernels loops annotation: C and C++ Frederik Harwath
2021-12-15 15:54 ` [PATCH 02/40] Add -fno-openacc-kernels-annotate-loops option to more testcases Frederik Harwath
2021-12-15 15:54 ` [PATCH 03/40] Kernels loops annotation: Fortran Frederik Harwath
2021-12-15 15:54 ` [PATCH 04/40] Additional Fortran testsuite fixes for kernels loops annotation pass Frederik Harwath
2021-12-15 15:54 ` [PATCH 05/40] Fix bug in processing of array dimensions in data clauses Frederik Harwath
2021-12-15 15:54 ` [PATCH 06/40] Add a "combined" flag for "acc kernels loop" etc directives Frederik Harwath
2021-12-15 15:54 ` [PATCH 07/40] Annotate inner loops in "acc kernels loop" directives (C/C++) Frederik Harwath
2021-12-15 15:54 ` [PATCH 08/40] Annotate inner loops in "acc kernels loop" directives (Fortran) Frederik Harwath
2021-12-15 15:54 ` [PATCH 09/40] Permit calls to builtins and intrinsics in kernels loops Frederik Harwath
2021-12-15 15:54 ` [PATCH 10/40] Fix patterns in Fortran tests for kernels loop annotation Frederik Harwath
2021-12-15 15:54 ` [PATCH 11/40] Clean up loop variable extraction in OpenACC " Frederik Harwath
2021-12-15 15:54 ` [PATCH 12/40] Relax some restrictions on the loop bound in " Frederik Harwath
2021-12-15 15:54 ` [PATCH 13/40] Fortran: Delinearize array accesses Frederik Harwath
2021-12-15 15:54 ` [PATCH 14/40] openacc: Move pass_oacc_device_lower after pass_graphite Frederik Harwath
2021-12-15 15:54 ` [PATCH 15/40] graphite: Extend SCoP detection dump output Frederik Harwath
2022-05-16 12:49 ` Tobias Burnus
2022-05-17 8:21 ` Richard Biener
2022-05-18 12:19 ` Harwath, Frederik
2022-05-18 12:21 ` Richard Biener
2021-12-15 15:54 ` [PATCH 16/40] graphite: Rename isl_id_for_ssa_name Frederik Harwath
2022-05-16 12:49 ` Tobias Burnus
2022-05-17 8:22 ` Richard Biener
2021-12-15 15:54 ` [PATCH 17/40] graphite: Fix minor mistakes in comments Frederik Harwath
2022-05-16 12:49 ` Tobias Burnus
2022-05-17 8:22 ` Richard Biener
2021-12-15 15:54 ` [PATCH 18/40] Move compute_alias_check_pairs to tree-data-ref.c Frederik Harwath
2021-12-15 15:54 ` [PATCH 19/40] graphite: Add runtime alias checking Frederik Harwath
2021-12-15 15:54 ` [PATCH 20/40] openacc: Use Graphite for dependence analysis in "kernels" regions Frederik Harwath
2021-12-15 15:54 ` [PATCH 21/40] openacc: Add "can_be_parallel" flag info to "graph" dumps Frederik Harwath
2021-12-15 15:54 ` [PATCH 22/40] openacc: Remove unused partitioning in "kernels" regions Frederik Harwath
2021-12-15 15:54 ` [PATCH 23/40] Add function for printing a single OMP_CLAUSE Frederik Harwath
2021-12-15 15:54 ` [PATCH 24/40] openacc: Add data optimization pass Frederik Harwath
2021-12-15 15:54 ` [PATCH 25/40] openacc: Add runtime alias checking for OpenACC kernels Frederik Harwath
2021-12-15 15:54 ` [PATCH 26/40] openacc: Warn about "independent" "kernels" loops with data-dependences Frederik Harwath
2021-12-15 15:54 ` [PATCH 27/40] openacc: Handle internal function calls in pass_lim Frederik Harwath
2021-12-15 15:54 ` [PATCH 28/40] openacc: Disable pass_pre on outlined functions analyzed by Graphite Frederik Harwath
2021-12-15 15:54 ` [PATCH 29/40] graphite: Tune parameters for OpenACC use Frederik Harwath
2021-12-15 15:54 ` [PATCH 30/40] graphite: Adjust scop loop-nest choice Frederik Harwath
2021-12-15 15:54 ` [PATCH 31/40] graphite: Accept loops without data references Frederik Harwath
2021-12-15 15:54 ` [PATCH 32/40] Reference reduction localization Frederik Harwath
2021-12-15 15:54 ` [PATCH 33/40] Fix tree check failure with " Frederik Harwath
2021-12-15 15:54 ` [PATCH 34/40] Use more appropriate var in localize_reductions call Frederik Harwath
2021-12-15 15:54 ` Frederik Harwath [this message]
2021-12-15 15:54 ` [PATCH 36/40] openacc: Enable reduction variable localization for "kernels" Frederik Harwath
2021-12-15 15:54 ` [PATCH 37/40] Fix for is_gimple_reg vars to 'data kernels' Frederik Harwath
2021-12-15 15:54 ` [PATCH 38/40] openacc: fix privatization of by-reference arrays Frederik Harwath
2021-12-15 15:54 ` [PATCH 39/40] openacc: Check type for references in reduction lowering Frederik Harwath
2021-12-16 12:00 ` [PATCH 40/40] openacc: Adjust testsuite to new "kernels" handling Frederik Harwath
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20211215155447.19379-36-frederik@codesourcery.com \
--to=frederik@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=julian@codesourcery.com \
--cc=thomas@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
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).