* [PATCH] [og9] Handle references in OpenACC "private" clauses
@ 2019-09-20 21:17 Julian Brown
2022-01-17 8:01 ` Test cases for references in OpenACC 'private' clauses Thomas Schwinge
0 siblings, 1 reply; 3+ messages in thread
From: Julian Brown @ 2019-09-20 21:17 UTC (permalink / raw)
To: gcc-patches; +Cc: Andrew Stubbs, Thomas Schwinge
This patch rewrites reference-type variables appearing in OpenACC
"private" clauses in a similar way to how such variables are handled in
reduction clauses. Otherwise, the mechanism used to privatize reference
variables is currently ill-suited to the worker-partitioning mechanism
used for AMD GCN, and each worker ends up accessing worker 0's copy of
those reference variables via broadcast pointers. Rewriting reference
variables to non-reference-type scalars sidesteps that problem.
This is intended as a somewhat temporary solution: it works for the
newly-included tests, but is not very elegant.
Tested with offloading to AMD GCN. I will apply to the
openacc-gcc-9-branch shortly.
Cheers,
Julian
ChangeLog
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/ChangeLog.openacc | 5 ++
gcc/gimplify.c | 15 ++++
libgomp/ChangeLog.openacc | 6 ++
.../libgomp.oacc-c++/privatized-ref-2.C | 64 +++++++++++++++++
.../libgomp.oacc-c++/privatized-ref-3.C | 64 +++++++++++++++++
.../libgomp.oacc-fortran/privatized-ref-1.f95 | 71 +++++++++++++++++++
6 files changed, 225 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/ChangeLog.openacc b/gcc/ChangeLog.openacc
index fe584959153..523b6eb1d74 100644
--- a/gcc/ChangeLog.openacc
+++ b/gcc/ChangeLog.openacc
@@ -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 d16611d3617..d95ad5d4baa 100644
--- 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.openacc b/libgomp/ChangeLog.openacc
index 7813760e642..d9d1c353e31 100644
--- a/libgomp/ChangeLog.openacc
+++ b/libgomp/ChangeLog.openacc
@@ -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
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
--
2.22.0
^ permalink raw reply [flat|nested] 3+ messages in thread
* Test cases for references in OpenACC 'private' clauses
2019-09-20 21:17 [PATCH] [og9] Handle references in OpenACC "private" clauses Julian Brown
@ 2022-01-17 8:01 ` Thomas Schwinge
2022-01-17 8:04 ` Extend test cases for references in OpenACC 'private' clauses (was: Test cases for references in OpenACC 'private' clauses) Thomas Schwinge
0 siblings, 1 reply; 3+ messages in thread
From: Thomas Schwinge @ 2022-01-17 8:01 UTC (permalink / raw)
To: gcc-patches; +Cc: Julian Brown
[-- Attachment #1: Type: text/plain, Size: 1373 bytes --]
Hi!
On 2019-09-20T14:17:33-0700, Julian Brown <julian@codesourcery.com> wrote:
> This patch ['Handle references in OpenACC "private" clauses'] [...]
..., and its prerequisite changes, in particular...
> [...] is intended as a somewhat temporary solution: it works for the
> newly-included tests, but is not very elegant.
..., and breaks other things, as discussed internally a while ago.
This will have to be done differently/analyzed in more detail.
The test cases however, amend with:
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
... etc., and 'libgomp.oacc-c++/privatized-ref-3.C',
'libgomp.oacc-fortran/privatized-ref-1.f95' with:
+/*TODO
+ { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */
... I've now pushed to master branch in
commit fbb438808e9b53a6e6b179a5787d609443acaad6
"Test cases for references in OpenACC 'private' clauses", see attached.
Grüße
Thomas
-----------------
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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Test-cases-for-references-in-OpenACC-private-clauses.patch --]
[-- Type: text/x-diff, Size: 6177 bytes --]
From fbb438808e9b53a6e6b179a5787d609443acaad6 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Fri, 20 Sep 2019 13:53:10 -0700
Subject: [PATCH] Test cases for references in OpenACC 'private' clauses
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.
Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
---
.../libgomp.oacc-c++/privatized-ref-2.C | 66 ++++++++++++++++
.../libgomp.oacc-c++/privatized-ref-3.C | 69 +++++++++++++++++
.../libgomp.oacc-fortran/privatized-ref-1.f95 | 76 +++++++++++++++++++
3 files changed, 211 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/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..7091091cac2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
@@ -0,0 +1,66 @@
+/* { 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)
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+ {
+ 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)
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+ {
+ 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..478876e3596
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C
@@ -0,0 +1,69 @@
+/* { dg-do run } */
+
+/*TODO
+ { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */
+
+#include <stdlib.h>
+
+void workers (void)
+{
+ double res[65536];
+ int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+ {
+ 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)
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+ {
+ 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..bb0910b1006
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95
@@ -0,0 +1,76 @@
+! { dg-do run }
+
+!TODO
+! { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } }
+
+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 } .-1 }
+
+ !$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 } .-1 }
+
+ !$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.34.1
^ permalink raw reply [flat|nested] 3+ messages in thread
* Extend test cases for references in OpenACC 'private' clauses (was: Test cases for references in OpenACC 'private' clauses)
2022-01-17 8:01 ` Test cases for references in OpenACC 'private' clauses Thomas Schwinge
@ 2022-01-17 8:04 ` Thomas Schwinge
0 siblings, 0 replies; 3+ messages in thread
From: Thomas Schwinge @ 2022-01-17 8:04 UTC (permalink / raw)
To: gcc-patches; +Cc: Julian Brown
[-- Attachment #1: Type: text/plain, Size: 655 bytes --]
Hi!
On 2022-01-17T09:01:54+0100, I wrote:
> I've now pushed to master branch in
> commit fbb438808e9b53a6e6b179a5787d609443acaad6
> "Test cases for references in OpenACC 'private' clauses", see attached.
... extended by commit b75aab194e3fe40b594d9a70eb7068dc9950bcf0
"Extend test cases for references in OpenACC 'private' clauses",
see attached.
Grüße
Thomas
-----------------
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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Extend-test-cases-for-references-in-OpenACC-private-.patch --]
[-- Type: text/x-diff, Size: 30731 bytes --]
From b75aab194e3fe40b594d9a70eb7068dc9950bcf0 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 24 Aug 2021 18:33:04 +0200
Subject: [PATCH] Extend test cases for references in OpenACC 'private' clauses
libgomp/
* testsuite/libgomp.oacc-c++/privatized-ref-2.C: Extend.
* testsuite/libgomp.oacc-c++/privatized-ref-3.C: Likewise.
* testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: Likewise.
---
.../libgomp.oacc-c++/privatized-ref-2.C | 128 +++++++++++++-
.../libgomp.oacc-c++/privatized-ref-3.C | 159 +++++++++++++++++-
.../libgomp.oacc-fortran/privatized-ref-1.f95 | 99 ++++++++++-
3 files changed, 364 insertions(+), 22 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
index 7091091cac2..520016ab59d 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C
@@ -1,20 +1,87 @@
/* { dg-do run } */
+/* { dg-additional-options "-fopt-info-note-omp" }
+ { dg-additional-options "-foffload=-fopt-info-note-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+
+/* { dg-additional-options "-Wuninitialized" } */
+
+/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+ passed to 'incr' may be unset, and in that case, it will be set to [...]",
+ so to maintain compatibility with earlier Tcl releases, we manually
+ initialize counter variables:
+ { dg-line l_dummy[variable c_compute 0 c_loop 0] }
+ { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+ "WARNING: dg-line var l_dummy defined, but not used". */
+
#include <stdlib.h>
+void gangs (void)
+{
+ double res[65536];
+ int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ {
+ int i, j;
+#pragma acc loop collapse(2) gang /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+ But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
+ No longer having address taken: tmpvar
+ Now a gimple register: tmpvar
+ However, 'tmpvar' remains in the candidate set:
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
+ Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_loop$c_loop }
+ For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_loop$c_loop }
+ { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_loop$c_loop }
+ */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ for (i = 0; i < 256; i++)
+ {
+ for (j = 0; j < 256; j++)
+ {
+ int tmpvar;
+ int &tmpref = tmpvar;
+ tmpref = (i * 256 + j) * 97;
+ res[i * 256 + j] = tmpref;
+ }
+ }
+ }
+
+ for (i = 0; i < 65536; i++)
+ if (res[i] != i * 97)
+ abort ();
+}
+
void workers (void)
{
double res[65536];
int i;
-#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
- /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
-#pragma acc loop gang
+#pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
-#pragma acc loop worker
+#pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'worker'} "" { target *-*-* } l_loop$c_loop }
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'worker'} "TODO" { target { ! openacc_host_selected } xfail *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (j = 0; j < 256; j++)
{
int tmpvar;
@@ -35,14 +102,22 @@ void vectors (void)
double res[65536];
int i;
-#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
- /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
-#pragma acc loop gang worker
+#pragma acc loop gang worker /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
-#pragma acc loop vector
+#pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop }
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'vector'} "TODO" { target { ! openacc_host_selected } } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (j = 0; j < 256; j++)
{
int tmpvar;
@@ -58,9 +133,46 @@ void vectors (void)
abort ();
}
+void gangs_workers_vectors (void)
+{
+ double res[65536];
+ int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
+ {
+ int i, j;
+#pragma acc loop collapse(2) gang worker vector /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop }
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'vector'} "TODO" { target { ! openacc_host_selected } } l_loop$c_loop } */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ for (i = 0; i < 256; i++)
+ {
+ for (j = 0; j < 256; j++)
+ {
+ int tmpvar;
+ int &tmpref = tmpvar;
+ tmpref = (i * 256 + j) * 103;
+ res[i * 256 + j] = tmpref;
+ }
+ }
+ }
+
+ for (i = 0; i < 65536; i++)
+ if (res[i] != i * 103)
+ abort ();
+}
+
int main (int argc, char *argv[])
{
+ gangs ();
workers ();
vectors ();
+ gangs_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
index 478876e3596..cb7085a01e8 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C
@@ -1,25 +1,103 @@
/* { dg-do run } */
+/* { dg-additional-options "-fopt-info-note-omp" }
+ { dg-additional-options "-foffload=-fopt-info-note-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+
+/* { dg-additional-options "-Wuninitialized" } */
+
/*TODO
{ dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */
+/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+ passed to 'incr' may be unset, and in that case, it will be set to [...]",
+ so to maintain compatibility with earlier Tcl releases, we manually
+ initialize counter variables:
+ { dg-line l_dummy[variable c_compute 0 c_loop 0] }
+ { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+ "WARNING: dg-line var l_dummy defined, but not used". */
+
#include <stdlib.h>
+void gangs (void)
+{
+ double res[65536];
+ int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+ But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
+ No longer having address taken: tmpvar
+ Now a gimple register: tmpvar
+ However, 'tmpvar' remains in the candidate set:
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
+ Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
+ For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
+ { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
+ */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ {
+ int i, j;
+ int tmpvar;
+ int &tmpref = tmpvar;
+#pragma acc loop collapse(2) gang private(tmpref) /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ for (i = 0; i < 256; i++)
+ {
+ for (j = 0; j < 256; j++)
+ {
+ tmpref = (i * 256 + j) * 97;
+ res[i * 256 + j] = tmpref;
+ }
+ }
+ }
+
+ for (i = 0; i < 65536; i++)
+ if (res[i] != i * 97)
+ abort ();
+}
+
void workers (void)
{
double res[65536];
int i;
-#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
- /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+ But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
+ No longer having address taken: tmpvar
+ Now a gimple register: tmpvar
+ However, 'tmpvar' remains in the candidate set:
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
+ Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
+ For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
+ { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
+ */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
int tmpvar;
int &tmpref = tmpvar;
-#pragma acc loop gang
+#pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
-#pragma acc loop worker private(tmpref)
+#pragma acc loop worker private(tmpref) /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (j = 0; j < 256; j++)
{
tmpref = (i * 256 + j) * 99;
@@ -38,16 +116,34 @@ void vectors (void)
double res[65536];
int i;
-#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
- /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+ But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
+ No longer having address taken: tmpvar
+ Now a gimple register: tmpvar
+ However, 'tmpvar' remains in the candidate set:
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
+ Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
+ For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
+ { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
+ */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
int tmpvar;
int &tmpref = tmpvar;
-#pragma acc loop gang worker
+#pragma acc loop gang worker /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
-#pragma acc loop vector private(tmpref)
+#pragma acc loop vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (j = 0; j < 256; j++)
{
tmpref = (i * 256 + j) * 101;
@@ -61,9 +157,56 @@ void vectors (void)
abort ();
}
+void gangs_workers_vectors (void)
+{
+ double res[65536];
+ int i;
+
+#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
+ /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+ But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
+ No longer having address taken: tmpvar
+ Now a gimple register: tmpvar
+ However, 'tmpvar' remains in the candidate set:
+ { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
+ Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
+ For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
+ { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
+ { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
+ */
+ /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+ /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
+ {
+ int i, j;
+ int tmpvar;
+ int &tmpref = tmpvar;
+#pragma acc loop collapse(2) gang worker vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
+ /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+ for (i = 0; i < 256; i++)
+ {
+ for (j = 0; j < 256; j++)
+ {
+ tmpref = (i * 256 + j) * 103;
+ res[i * 256 + j] = tmpref;
+ }
+ }
+ }
+
+ for (i = 0; i < 65536; i++)
+ if (res[i] != i * 103)
+ abort ();
+}
+
int main (int argc, char *argv[])
{
+ gangs ();
workers ();
vectors ();
+ gangs_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
index bb0910b1006..a8230561fc9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95
@@ -1,8 +1,24 @@
! { dg-do run }
+! { dg-additional-options "-fopt-info-note-omp" }
+! { dg-additional-options "-foffload=-fopt-info-note-omp" }
+
+! { dg-additional-options "--param=openacc-privatization=noisy" }
+! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+
+! { dg-additional-options "-Wuninitialized" }
+
!TODO
! { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } }
+! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+! passed to 'incr' may be unset, and in that case, it will be set to [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c_loop 0] }
+! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used". */
+
program main
implicit none
integer :: myint
@@ -11,12 +27,22 @@ program main
res(:) = 0.0
+ myint = 3
+ call gangs(myint, res)
+
+ do i=1,65536
+ tmp = i * 97
+ if (res(i) .ne. tmp) stop 1
+ end do
+
+ res(:) = 0.0
+
myint = 5
call workers(myint, res)
do i=1,65536
tmp = i * 99
- if (res(i) .ne. tmp) stop 1
+ if (res(i) .ne. tmp) stop 2
end do
res(:) = 0.0
@@ -26,11 +52,43 @@ program main
do i=1,65536
tmp = i * 101
- if (res(i) .ne. tmp) stop 2
+ if (res(i) .ne. tmp) stop 3
+ end do
+
+ res(:) = 0.0
+
+ myint = 9
+ call gangs_workers_vectors(myint, res)
+
+ do i=1,65536
+ tmp = i * 103
+ if (res(i) .ne. tmp) stop 4
end do
contains
+ subroutine gangs(t1, res)
+ implicit none
+ integer :: t1
+ integer :: i, j
+ real, intent(out) :: res(:)
+
+ !$acc parallel copyout(res) num_gangs(64)
+
+ !$acc loop collapse(2) gang private(t1) ! { dg-line l_loop[incr c_loop] }
+ ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+ ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+ ! { dg-note {variable 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
+ do i=0,255
+ do j=1,256
+ t1 = (i * 256 + j) * 97
+ res(i * 256 + j) = t1
+ end do
+ end do
+
+ !$acc end parallel
+ end subroutine gangs
+
subroutine workers(t1, res)
implicit none
integer :: t1
@@ -40,9 +98,12 @@ contains
!$acc parallel copyout(res) num_gangs(64) num_workers(64)
! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 }
- !$acc loop gang
+ !$acc loop gang ! { dg-line l_loop[incr c_loop] }
+ ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
do i=0,255
- !$acc loop worker private(t1)
+ !$acc loop worker private(t1) ! { dg-line l_loop[incr c_loop] }
+ ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+ ! { dg-note {variable 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
do j=1,256
t1 = (i * 256 + j) * 99
res(i * 256 + j) = t1
@@ -61,9 +122,12 @@ contains
!$acc parallel copyout(res) num_gangs(64) num_workers(64)
! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 }
- !$acc loop gang worker
+ !$acc loop gang worker ! { dg-line l_loop[incr c_loop] }
+ ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
do i=0,255
- !$acc loop vector private(t1)
+ !$acc loop vector private(t1) ! { dg-line l_loop[incr c_loop] }
+ ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+ ! { dg-note {variable 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
do j=1,256
t1 = (i * 256 + j) * 101
res(i * 256 + j) = t1
@@ -73,4 +137,27 @@ contains
!$acc end parallel
end subroutine vectors
+ subroutine gangs_workers_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 } .-1 }
+
+ !$acc loop collapse(2) gang worker vector private(t1) ! { dg-line l_loop[incr c_loop] }
+ ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+ ! { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+ ! { dg-note {variable 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
+ do i=0,255
+ do j=1,256
+ t1 = (i * 256 + j) * 103
+ res(i * 256 + j) = t1
+ end do
+ end do
+
+ !$acc end parallel
+ end subroutine gangs_workers_vectors
+
end program main
--
2.34.1
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2022-01-17 8:04 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2019-09-20 21:17 [PATCH] [og9] Handle references in OpenACC "private" clauses Julian Brown
2022-01-17 8:01 ` Test cases for references in OpenACC 'private' clauses Thomas Schwinge
2022-01-17 8:04 ` Extend test cases for references in OpenACC 'private' clauses (was: Test cases for references in OpenACC 'private' clauses) Thomas Schwinge
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).