* [patch,gomp4] make fortran loop variables implicitly private in openacc
@ 2014-08-11 23:55 Cesar Philippidis
2014-08-13 3:55 ` Cesar Philippidis
` (2 more replies)
0 siblings, 3 replies; 12+ messages in thread
From: Cesar Philippidis @ 2014-08-11 23:55 UTC (permalink / raw)
To: fortran, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 431 bytes --]
According to section 2.6.1 in the openacc spec, fortran loop variables
should be implicitly private like in openmp. This patch does just so.
Also, while working on this patch, I noticed that I made the check for
variables appearing in multiple openacc clauses too strict. A private
variable may also appear inside a reduction clause. I've also included a
fix for this in this patch.
Is this OK for gomp-4_0-branch?
Thanks,
Cesar
[-- Attachment #2: gfortran-acc-do-private.diff --]
[-- Type: text/x-patch, Size: 4637 bytes --]
2014-08-11 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* openmp.c (oacc_compatible_clauses): New function.
(resolve_omp_clauses): Use it.
(oacc_current_ctx): Move it near omp_current_ctx.
(gfc_resolve_do_iterator): Handle OpenACC index variables.
(gfc_resolve_oacc_blocks): Initialize ctx.share_clauses and
ctx.private_iterators.
gcc/testsuite/
* gfortran.dg/goacc/private-1.f95: New test.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 91e00c4..2c91597 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2713,6 +2713,29 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
return copy;
}
+/* Returns true if clause in list 'list' is compatible with any of
+ of the clauses in lists [0..list-1]. E.g., a reduction variable may
+ appear in both reduction and private clauses, so this function
+ will return true in this case. */
+
+static bool
+oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
+ gfc_symbol *sym, bool openacc)
+{
+ gfc_omp_namelist *n;
+
+ if (!openacc)
+ return false;
+
+ if (list != OMP_LIST_REDUCTION)
+ return false;
+
+ for (n = clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
+ if (n->sym == sym)
+ return true;
+
+ return false;
+}
/* OpenMP directive resolving routines. */
@@ -2826,7 +2849,8 @@ resolve_omp_clauses (gfc_code *code, locus *where,
&& list != OMP_LIST_TO)
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark)
+ if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
+ n->sym, openacc))
gfc_error ("Symbol '%s' present on multiple clauses at %L",
n->sym->name, where);
else
@@ -3791,6 +3815,9 @@ struct omp_context
static gfc_code *omp_current_do_code;
static int omp_current_do_collapse;
+typedef struct omp_context oacc_context;
+oacc_context *oacc_current_ctx;
+
void
gfc_resolve_omp_do_blocks (gfc_code *code, gfc_namespace *ns)
{
@@ -3906,6 +3933,8 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym)
{
int i = omp_current_do_collapse;
gfc_code *c = omp_current_do_code;
+ bool openacc = omp_current_ctx == NULL;
+ omp_context *current_ctx = openacc ? oacc_current_ctx : omp_current_ctx;
if (sym->attr.threadprivate)
return;
@@ -3922,15 +3951,15 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym)
c = c->block->next;
}
- if (omp_current_ctx == NULL)
+ if (current_ctx == NULL)
return;
- if (pointer_set_contains (omp_current_ctx->sharing_clauses, sym))
+ if (!openacc && pointer_set_contains (current_ctx->sharing_clauses, sym))
return;
- if (! pointer_set_insert (omp_current_ctx->private_iterators, sym))
+ if (! pointer_set_insert (current_ctx->private_iterators, sym))
{
- gfc_omp_clauses *omp_clauses = omp_current_ctx->code->ext.omp_clauses;
+ gfc_omp_clauses *omp_clauses = current_ctx->code->ext.omp_clauses;
gfc_omp_namelist *p;
p = gfc_get_omp_namelist ();
@@ -4106,9 +4135,6 @@ resolve_omp_do (gfc_code *code)
}
}
-typedef struct omp_context oacc_context;
-oacc_context *oacc_current_ctx;
-
static bool
oacc_is_parallel (gfc_code *code)
{
@@ -4424,6 +4450,8 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
resolve_oacc_loop_blocks (code);
ctx.code = code;
+ ctx.sharing_clauses = NULL;
+ ctx.private_iterators = pointer_set_create ();
ctx.previous = oacc_current_ctx;
oacc_current_ctx = &ctx;
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
new file mode 100644
index 0000000..4eaec4f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k
+ logical :: l
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+end program test
+! { dg-prune-output "unimplemented" }
+! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [patch,gomp4] make fortran loop variables implicitly private in openacc
2014-08-11 23:55 [patch,gomp4] make fortran loop variables implicitly private in openacc Cesar Philippidis
@ 2014-08-13 3:55 ` Cesar Philippidis
2014-08-13 20:41 ` Tobias Burnus
2017-01-27 15:03 ` [gomp4] don't error on implicitly private induction variables in gfortran Cesar Philippidis
2 siblings, 0 replies; 12+ messages in thread
From: Cesar Philippidis @ 2014-08-13 3:55 UTC (permalink / raw)
To: fortran, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1148 bytes --]
On 08/11/2014 04:55 PM, Cesar Philippidis wrote:
> According to section 2.6.1 in the openacc spec, fortran loop variables
> should be implicitly private like in openmp. This patch does just so.
> Also, while working on this patch, I noticed that I made the check for
> variables appearing in multiple openacc clauses too strict. A private
> variable may also appear inside a reduction clause. I've also included a
> fix for this in this patch.
>
> Is this OK for gomp-4_0-branch?
I've updated this patch to properly handle loop nests inside openacc
data blocks. In the original patch, something like this
!$acc data copy(A)
do z = 1, 100
!$acc parallel
!$acc loop
do j=1,m-2
end do
!$acc end parallel
end do
would result in the loop variable 'z' being implicitly treated as
private. This occurs because gfc_resolve_oacc_blocks updates
oacc_current_ctx for both loop and data blocks. Originally
omp_current_ctx was only associated with do blocks, so
gfc_resolve_do_iterator didn't expect non-loop ctx's. This revised patch
makes gfc_resolve_do_iterator aware potential data blocks.
Is this OK for gomp-4_0-branch?
Thanks,
Cesar
[-- Attachment #2: gfortran-acc-do-private-b.diff --]
[-- Type: text/x-patch, Size: 5771 bytes --]
2014-08-11 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* openmp.c (oacc_compatible_clauses): New function.
(resolve_omp_clauses): Use it.
(oacc_current_ctx): Move it near omp_current_ctx.
(gfc_resolve_do_iterator): Handle OpenACC index variables.
(gfc_resolve_oacc_blocks): Initialize ctx.share_clauses and
ctx.private_iterators.
gcc/testsuite/
* gfortran.dg/goacc/private-1.f95: New test.
* gfortran.dg/goacc/private-2.f95: New test.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 91e00c4..4bbbf2f 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2713,6 +2713,29 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
return copy;
}
+/* Returns true if clause in list 'list' is compatible with any of
+ of the clauses in lists [0..list-1]. E.g., a reduction variable may
+ appear in both reduction and private clauses, so this function
+ will return true in this case. */
+
+static bool
+oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
+ gfc_symbol *sym, bool openacc)
+{
+ gfc_omp_namelist *n;
+
+ if (!openacc)
+ return false;
+
+ if (list != OMP_LIST_REDUCTION)
+ return false;
+
+ for (n = clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
+ if (n->sym == sym)
+ return true;
+
+ return false;
+}
/* OpenMP directive resolving routines. */
@@ -2826,7 +2849,8 @@ resolve_omp_clauses (gfc_code *code, locus *where,
&& list != OMP_LIST_TO)
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark)
+ if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
+ n->sym, openacc))
gfc_error ("Symbol '%s' present on multiple clauses at %L",
n->sym->name, where);
else
@@ -3791,6 +3815,9 @@ struct omp_context
static gfc_code *omp_current_do_code;
static int omp_current_do_collapse;
+typedef struct omp_context oacc_context;
+oacc_context *oacc_current_ctx;
+
void
gfc_resolve_omp_do_blocks (gfc_code *code, gfc_namespace *ns)
{
@@ -3906,6 +3933,8 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym)
{
int i = omp_current_do_collapse;
gfc_code *c = omp_current_do_code;
+ bool openacc = omp_current_ctx == NULL;
+ omp_context *current_ctx = openacc ? oacc_current_ctx : omp_current_ctx;
if (sym->attr.threadprivate)
return;
@@ -3922,15 +3951,19 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym)
c = c->block->next;
}
- if (omp_current_ctx == NULL)
+ if (current_ctx == NULL)
+ return;
+
+ /* An openacc context may represent a data clause. Abort if so. */
+ if (openacc && !oacc_is_loop (current_ctx->code))
return;
- if (pointer_set_contains (omp_current_ctx->sharing_clauses, sym))
+ if (!openacc && pointer_set_contains (current_ctx->sharing_clauses, sym))
return;
- if (! pointer_set_insert (omp_current_ctx->private_iterators, sym))
+ if (! pointer_set_insert (current_ctx->private_iterators, sym))
{
- gfc_omp_clauses *omp_clauses = omp_current_ctx->code->ext.omp_clauses;
+ gfc_omp_clauses *omp_clauses = current_ctx->code->ext.omp_clauses;
gfc_omp_namelist *p;
p = gfc_get_omp_namelist ();
@@ -4106,9 +4139,6 @@ resolve_omp_do (gfc_code *code)
}
}
-typedef struct omp_context oacc_context;
-oacc_context *oacc_current_ctx;
-
static bool
oacc_is_parallel (gfc_code *code)
{
@@ -4424,11 +4454,14 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
resolve_oacc_loop_blocks (code);
ctx.code = code;
+ ctx.sharing_clauses = NULL;
+ ctx.private_iterators = pointer_set_create ();
ctx.previous = oacc_current_ctx;
oacc_current_ctx = &ctx;
gfc_resolve_blocks (code->block, ns);
+ pointer_set_destroy (ctx.private_iterators);
oacc_current_ctx = ctx.previous;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
new file mode 100644
index 0000000..5aeee3b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
@@ -0,0 +1,38 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+end program test
+! { dg-prune-output "unimplemented" }
+! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-2.f95 b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
new file mode 100644
index 0000000..4b038f2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
@@ -0,0 +1,39 @@
+! { dg-do compile }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k, a(10)
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc data copy(a)
+
+ if(mod(1,10) .eq. 0) write(*,'(i5)') i
+
+ do i = 1, 100
+ !$acc parallel
+ !$acc loop
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ !$acc end parallel
+ end do
+
+ !$acc end data
+
+end program test
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [patch,gomp4] make fortran loop variables implicitly private in openacc
2014-08-11 23:55 [patch,gomp4] make fortran loop variables implicitly private in openacc Cesar Philippidis
2014-08-13 3:55 ` Cesar Philippidis
@ 2014-08-13 20:41 ` Tobias Burnus
2014-08-14 11:15 ` Thomas Schwinge
2014-08-22 15:08 ` Cesar Philippidis
2017-01-27 15:03 ` [gomp4] don't error on implicitly private induction variables in gfortran Cesar Philippidis
2 siblings, 2 replies; 12+ messages in thread
From: Tobias Burnus @ 2014-08-13 20:41 UTC (permalink / raw)
To: Cesar Philippidis, fortran, gcc-patches
Cesar Philippidis wrote:
> According to section 2.6.1 in the openacc spec, fortran loop variables
> should be implicitly private like in openmp. This patch does just so.
Makes sense. Looking at the patch, I wonder whether the context is
properly handled when mixing OpenMP and OpenACC. I have the feeling that
one then might end up using the OpenACC context when one should use the
OpenMP context. However, I have not fully followed the program flow. For
instance, would something like
!$oacc parallel
!$omp simd private(i) reduction(+:c)
do i = 1, n
...
end do
be properly handled?
> Also, while working on this patch, I noticed that I made the check for
> variables appearing in multiple openacc clauses too strict. A private
> variable may also appear inside a reduction clause. I've also included a
> fix for this in this patch.
In OpenMP, one has (OMP 4.0, 2.14.3): "A list item that specifies a
given variable may not appear in more than one clause on the same
directive, except that a variable may be specified in both firstprivate
and lastprivate clauses."
And in 2.14.3.3, OpenMP has: "List items that appear in a private,
firstprivate, or reduction clause in a parallel construct may also
appear in a private clause in an enclosed parallel,task, or worksharing,
or simd construct.
I tried to find it in something similar in OpenACC - but I failed. I
found in (OpenACC 2.0a, 2.7.11) a reference to reduction with private,
which implies that a reduction variable my be private, but I didn't find
much more. In particular, it is not clear to me whether it would permit
only those which are private implicitly or in an oacc parallel region or
also in an explicit private clause in the same directive. Can you point
me to the spec?
Additionally, I wonder whether you shouldn't also add a test case for
the reduction with private.
Tobias
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [patch,gomp4] make fortran loop variables implicitly private in openacc
2014-08-13 20:41 ` Tobias Burnus
@ 2014-08-14 11:15 ` Thomas Schwinge
2014-08-22 15:08 ` Cesar Philippidis
1 sibling, 0 replies; 12+ messages in thread
From: Thomas Schwinge @ 2014-08-14 11:15 UTC (permalink / raw)
To: Tobias Burnus, Cesar Philippidis; +Cc: fortran, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1038 bytes --]
Hi!
On Wed, 13 Aug 2014 22:41:47 +0200, Tobias Burnus <burnus@net-b.de> wrote:
> Cesar Philippidis wrote:
> > According to section 2.6.1 in the openacc spec, fortran loop variables
> > should be implicitly private like in openmp. This patch does just so.
>
> Makes sense. Looking at the patch, I wonder whether the context is
> properly handled when mixing OpenMP and OpenACC. I have the feeling that
> one then might end up using the OpenACC context when one should use the
> OpenMP context. However, I have not fully followed the program flow. For
> instance, would something like
>
> !$oacc parallel
> !$omp simd private(i) reduction(+:c)
> do i = 1, n
> ...
> end do
>
> be properly handled?
While that isn't something we're currently focussing on supporting, my
intuition tells me that there shouldn't be a separate oacc_current_ctx,
and omp_current_ctx should be used instead for OpenACC contexts, too.
(If you guys agree, don't let fixing that hold up the patch as posted.)
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [patch,gomp4] make fortran loop variables implicitly private in openacc
2014-08-13 20:41 ` Tobias Burnus
2014-08-14 11:15 ` Thomas Schwinge
@ 2014-08-22 15:08 ` Cesar Philippidis
2014-09-08 6:26 ` Tobias Burnus
1 sibling, 1 reply; 12+ messages in thread
From: Cesar Philippidis @ 2014-08-22 15:08 UTC (permalink / raw)
To: Tobias Burnus, fortran, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2705 bytes --]
On 08/13/2014 01:41 PM, Tobias Burnus wrote:
> Cesar Philippidis wrote:
>> According to section 2.6.1 in the openacc spec, fortran loop variables
>> should be implicitly private like in openmp. This patch does just so.
>
> Makes sense. Looking at the patch, I wonder whether the context is
> properly handled when mixing OpenMP and OpenACC. I have the feeling that
> one then might end up using the OpenACC context when one should use the
> OpenMP context. However, I have not fully followed the program flow. For
> instance, would something like
>
> !$oacc parallel
> !$omp simd private(i) reduction(+:c)
> do i = 1, n
> ...
> end do
>
> be properly handled?
These contexts are used to prevent situations like that. See
gfortran.dg/goacc/omp.f95 for examples. Also, Thomas mentioned that we
shouldn't need a separate oacc_context, so I've added an is_openmp field
to omp_context to indicate if omp_current_context represents an openmp
context (true) or openacc (false).
>> Also, while working on this patch, I noticed that I made the check for
>> variables appearing in multiple openacc clauses too strict. A private
>> variable may also appear inside a reduction clause. I've also included a
>> fix for this in this patch.
>
> In OpenMP, one has (OMP 4.0, 2.14.3): "A list item that specifies a
> given variable may not appear in more than one clause on the same
> directive, except that a variable may be specified in both firstprivate
> and lastprivate clauses."
> And in 2.14.3.3, OpenMP has: "List items that appear in a private,
> firstprivate, or reduction clause in a parallel construct may also
> appear in a private clause in an enclosed parallel,task, or worksharing,
> or simd construct.
>
> I tried to find it in something similar in OpenACC - but I failed. I
> found in (OpenACC 2.0a, 2.7.11) a reference to reduction with private,
> which implies that a reduction variable my be private, but I didn't find
> much more. In particular, it is not clear to me whether it would permit
> only those which are private implicitly or in an oacc parallel region or
> also in an explicit private clause in the same directive. Can you point
> me to the spec?
I'm not sure. I sent an email to the openacc technical mailing list, but
I haven't heard back from them.
> Additionally, I wonder whether you shouldn't also add a test case for
> the reduction with private.
I added one, see private-3.f95. Beware, part of that test is commented
out because our support for the private clause is incomplete. Thomas is
working on the private and firstprivate clauses, so he can uncomment
that portion of the test once the private clause is fully working.
Is this patch ok for gomp-4_0-branch?
Cesar
[-- Attachment #2: gfortran-acc-do-private-c.diff --]
[-- Type: text/x-patch, Size: 8709 bytes --]
2014-08-21 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* openmp.c (oacc_compatible_clauses): New function.
(resolve_omp_clauses): Use it.
(struct omp_context): Add is_openmp member.
(gfc_resolve_omp_parallel_blocks): Set is_openmp true.
(gfc_resolve_do_iterator): Scan for compatible clauses.
(typedef oacc_context): Remove.
(oacc_current_ctx): Remove. Use omp_current_ctx for both
OpenACC and OpenMP.
(resolve_oacc_directive_inside_omp_region): Replace
oacc_current_ctx with omp_current_ctx.
(resolve_omp_directive_inside_oacc_region): Likewise.
(resolve_oacc_nested_loops): Likewise.
(resolve_oacc_params_in_parallel): Likewise.
(resolve_oacc_loop_blocks): Likewise. Set is_openmp to false.
gcc/testsuite/
* gfortran.dg/goacc/private-1.f95: New test.
* gfortran.dg/goacc/private-2.f95: New test.
* gfortran.dg/goacc/private-3.f95: New test.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 91e00c4..e07a508 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2713,6 +2713,29 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
return copy;
}
+/* Returns true if clause in list 'list' is compatible with any of
+ of the clauses in lists [0..list-1]. E.g., a reduction variable may
+ appear in both reduction and private clauses, so this function
+ will return true in this case. */
+
+static bool
+oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
+ gfc_symbol *sym, bool openacc)
+{
+ gfc_omp_namelist *n;
+
+ if (!openacc)
+ return false;
+
+ if (list != OMP_LIST_REDUCTION)
+ return false;
+
+ for (n = clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
+ if (n->sym == sym)
+ return true;
+
+ return false;
+}
/* OpenMP directive resolving routines. */
@@ -2826,7 +2849,8 @@ resolve_omp_clauses (gfc_code *code, locus *where,
&& list != OMP_LIST_TO)
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark)
+ if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
+ n->sym, openacc))
gfc_error ("Symbol '%s' present on multiple clauses at %L",
n->sym->name, where);
else
@@ -3787,6 +3811,7 @@ struct omp_context
struct pointer_set_t *sharing_clauses;
struct pointer_set_t *private_iterators;
struct omp_context *previous;
+ bool is_openmp;
} *omp_current_ctx;
static gfc_code *omp_current_do_code;
static int omp_current_do_collapse;
@@ -3831,6 +3856,7 @@ gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns)
ctx.sharing_clauses = pointer_set_create ();
ctx.private_iterators = pointer_set_create ();
ctx.previous = omp_current_ctx;
+ ctx.is_openmp = true;
omp_current_ctx = &ctx;
for (list = 0; list < OMP_LIST_NUM; list++)
@@ -3925,7 +3951,12 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym)
if (omp_current_ctx == NULL)
return;
- if (pointer_set_contains (omp_current_ctx->sharing_clauses, sym))
+ /* An openacc context may represent a data clause. Abort if so. */
+ if (!omp_current_ctx->is_openmp && !oacc_is_loop (omp_current_ctx->code))
+ return;
+
+ if (omp_current_ctx->is_openmp
+ && pointer_set_contains (omp_current_ctx->sharing_clauses, sym))
return;
if (! pointer_set_insert (omp_current_ctx->private_iterators, sym))
@@ -4106,9 +4137,6 @@ resolve_omp_do (gfc_code *code)
}
}
-typedef struct omp_context oacc_context;
-oacc_context *oacc_current_ctx;
-
static bool
oacc_is_parallel (gfc_code *code)
{
@@ -4180,7 +4208,7 @@ switch (code->op)
static void
resolve_oacc_directive_inside_omp_region (gfc_code *code)
{
- if (omp_current_ctx != NULL)
+ if (omp_current_ctx != NULL && omp_current_ctx->is_openmp)
{
gfc_statement st = omp_code_to_statement (omp_current_ctx->code);
gfc_statement oacc_st = oacc_code_to_statement (code);
@@ -4193,9 +4221,9 @@ resolve_oacc_directive_inside_omp_region (gfc_code *code)
static void
resolve_omp_directive_inside_oacc_region (gfc_code *code)
{
- if (oacc_current_ctx != NULL)
+ if (omp_current_ctx != NULL && !omp_current_ctx->is_openmp)
{
- gfc_statement st = oacc_code_to_statement (oacc_current_ctx->code);
+ gfc_statement st = oacc_code_to_statement (omp_current_ctx->code);
gfc_statement omp_st = omp_code_to_statement (code);
gfc_error ("The %s directive cannot be specified within "
"a %s region at %L", gfc_ascii_statement (omp_st),
@@ -4282,12 +4310,12 @@ resolve_oacc_nested_loops (gfc_code *code, gfc_code* do_code, int collapse,
static void
resolve_oacc_params_in_parallel (gfc_code *code, const char *clause)
{
- oacc_context *c;
+ omp_context *c;
if (oacc_is_parallel (code))
gfc_error ("!$ACC LOOP %s in PARALLEL region doesn't allow "
"non-static arguments at %L", clause, &code->loc);
- for (c = oacc_current_ctx; c; c = c->previous)
+ for (c = omp_current_ctx; c; c = c->previous)
{
if (oacc_is_loop (c->code))
break;
@@ -4301,13 +4329,13 @@ resolve_oacc_params_in_parallel (gfc_code *code, const char *clause)
static void
resolve_oacc_loop_blocks (gfc_code *code)
{
- oacc_context *c;
+ omp_context *c;
if (!oacc_is_loop (code))
return;
if (code->op == EXEC_OACC_LOOP)
- for (c = oacc_current_ctx; c; c = c->previous)
+ for (c = omp_current_ctx; c; c = c->previous)
{
if (oacc_is_loop (c->code))
{
@@ -4419,17 +4447,21 @@ resolve_oacc_loop_blocks (gfc_code *code)
void
gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
{
- oacc_context ctx;
+ omp_context ctx;
resolve_oacc_loop_blocks (code);
ctx.code = code;
- ctx.previous = oacc_current_ctx;
- oacc_current_ctx = &ctx;
+ ctx.sharing_clauses = NULL;
+ ctx.private_iterators = pointer_set_create ();
+ ctx.previous = omp_current_ctx;
+ ctx.is_openmp = false;
+ omp_current_ctx = &ctx;
gfc_resolve_blocks (code->block, ns);
- oacc_current_ctx = ctx.previous;
+ pointer_set_destroy (ctx.private_iterators);
+ omp_current_ctx = ctx.previous;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
new file mode 100644
index 0000000..5aeee3b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
@@ -0,0 +1,38 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+end program test
+! { dg-prune-output "unimplemented" }
+! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-2.f95 b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
new file mode 100644
index 0000000..4b038f2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
@@ -0,0 +1,39 @@
+! { dg-do compile }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k, a(10)
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc data copy(a)
+
+ if(mod(1,10) .eq. 0) write(*,'(i5)') i
+
+ do i = 1, 100
+ !$acc parallel
+ !$acc loop
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ !$acc end parallel
+ end do
+
+ !$acc end data
+
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-3.f95 b/gcc/testsuite/gfortran.dg/goacc/private-3.f95
new file mode 100644
index 0000000..aa12a56
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-3.f95
@@ -0,0 +1,23 @@
+! { dg-do compile }
+
+! test for private variables in a reduction clause
+
+program test
+ implicit none
+ integer, parameter :: n = 100
+ integer :: i, k
+
+! FIXME: This causes an ICE in the gimplifier.
+! !$acc parallel private (k) reduction (+:k)
+! do i = 1, n
+! k = k + 1
+! end do
+! !$acc end parallel
+
+ !$acc parallel private (k)
+ !$acc loop reduction (+:k)
+ do i = 1, n
+ k = k + 1
+ end do
+ !$acc end parallel
+end program test
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [patch,gomp4] make fortran loop variables implicitly private in openacc
2014-08-22 15:08 ` Cesar Philippidis
@ 2014-09-08 6:26 ` Tobias Burnus
2014-09-09 1:25 ` Cesar Philippidis
0 siblings, 1 reply; 12+ messages in thread
From: Tobias Burnus @ 2014-09-08 6:26 UTC (permalink / raw)
To: Cesar Philippidis, fortran, gcc-patches
Dear Cesar,
sorry for the slow review.
On 22 August 2014 17:08, Cesar Philippidis wrote:
>> In OpenMP, one has (OMP 4.0, 2.14.3): "A list item that specifies a
>> given variable may not appear in more than one clause on the same
>> directive, except that a variable may be specified in both firstprivate
>> and lastprivate clauses."
>> And in 2.14.3.3, OpenMP has: "List items that appear in a private,
>> firstprivate, or reduction clause in a parallel construct may also
>> appear in a private clause in an enclosed parallel,task, or worksharing,
>> or simd construct.
>>
>> I tried to find it in something similar in OpenACC - but I failed. I
>> found in (OpenACC 2.0a, 2.7.11) a reference to reduction with private,
>> which implies that a reduction variable my be private, but I didn't find
>> much more. In particular, it is not clear to me whether it would permit
>> only those which are private implicitly or in an oacc parallel region or
>> also in an explicit private clause in the same directive. Can you point
>> me to the spec?
> I'm not sure. I sent an email to the openacc technical mailing list, but
> I haven't heard back from them.
Any new on this?
> Is this patch ok for gomp-4_0-branch?
Looks good to me.
Tobias
PS: Looking at
+ for (n = clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
I was stumbing a bit until I realized that OMP_LIST_PRIVATE is the first
element. I wonder whether one should have something like OMP_LIST_FIRST
= OMP_LIST_PRIVATE and OMP_LIST_LAST = OMP_LIST_NUM (or
OMP_LIST_FIRST_ENUM or ...) which avoids the implicit knowledge which
comes first in the enum.
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [patch,gomp4] make fortran loop variables implicitly private in openacc
2014-09-08 6:26 ` Tobias Burnus
@ 2014-09-09 1:25 ` Cesar Philippidis
0 siblings, 0 replies; 12+ messages in thread
From: Cesar Philippidis @ 2014-09-09 1:25 UTC (permalink / raw)
To: Tobias Burnus, fortran, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1907 bytes --]
On 09/07/2014 11:26 PM, Tobias Burnus wrote:
> sorry for the slow review.
No problem. I've been focusing on the subroutine clause lately.
> On 22 August 2014 17:08, Cesar Philippidis wrote:
>>> In OpenMP, one has (OMP 4.0, 2.14.3): "A list item that specifies a
>>> given variable may not appear in more than one clause on the same
>>> directive, except that a variable may be specified in both firstprivate
>>> and lastprivate clauses."
>>> And in 2.14.3.3, OpenMP has: "List items that appear in a private,
>>> firstprivate, or reduction clause in a parallel construct may also
>>> appear in a private clause in an enclosed parallel,task, or worksharing,
>>> or simd construct.
>>>
>>> I tried to find it in something similar in OpenACC - but I failed. I
>>> found in (OpenACC 2.0a, 2.7.11) a reference to reduction with private,
>>> which implies that a reduction variable my be private, but I didn't find
>>> much more. In particular, it is not clear to me whether it would permit
>>> only those which are private implicitly or in an oacc parallel region or
>>> also in an explicit private clause in the same directive. Can you point
>>> me to the spec?
>> I'm not sure. I sent an email to the openacc technical mailing list, but
>> I haven't heard back from them.
>
> Any new on this?
Not yet, unfortunately.
>> Is this patch ok for gomp-4_0-branch?
>
> Looks good to me.
>
> Tobias
>
> PS: Looking at
>
> + for (n = clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
>
> I was stumbing a bit until I realized that OMP_LIST_PRIVATE is the first
> element. I wonder whether one should have something like OMP_LIST_FIRST
> = OMP_LIST_PRIVATE and OMP_LIST_LAST = OMP_LIST_NUM (or
> OMP_LIST_FIRST_ENUM or ...) which avoids the implicit knowledge which
> comes first in the enum.
That's a good idea. I've made that change in the attached patch and
committed to gomp-4_0-branch.
Thanks,
Cesar
[-- Attachment #2: gfortran-acc-do-private-d.diff --]
[-- Type: text/x-patch, Size: 9352 bytes --]
2014-09-08 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* gfortran.h (enum OMP_LIST_FIRST, OMP_LIST_LAST): New
OMP enums.
* openmp.c (oacc_compatible_clauses): New function.
(resolve_omp_clauses): Use it.
(struct omp_context): Add is_openmp member.
(gfc_resolve_omp_parallel_blocks): Set is_openmp true.
(gfc_resolve_do_iterator): Scan for compatible clauses.
(typedef oacc_context): Remove.
(oacc_current_ctx): Remove. Use omp_current_ctx for both
OpenACC and OpenMP.
(resolve_oacc_directive_inside_omp_region): Replace
oacc_current_ctx with omp_current_ctx.
(resolve_omp_directive_inside_oacc_region): Likewise.
(resolve_oacc_nested_loops): Likewise.
(resolve_oacc_params_in_parallel): Likewise.
(resolve_oacc_loop_blocks): Likewise. Set is_openmp to false.
gcc/testsuite/
* gfortran.dg/goacc/private-1.f95: New test.
* gfortran.dg/goacc/private-2.f95: New test.
* gfortran.dg/goacc/private-3.f95: New test.
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 0cde668..abe4f05 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1143,7 +1143,8 @@ gfc_omp_namelist;
enum
{
- OMP_LIST_PRIVATE,
+ OMP_LIST_FIRST,
+ OMP_LIST_PRIVATE = OMP_LIST_FIRST,
OMP_LIST_FIRSTPRIVATE,
OMP_LIST_LASTPRIVATE,
OMP_LIST_COPYPRIVATE,
@@ -1166,7 +1167,8 @@ enum
OMP_LIST_HOST,
OMP_LIST_DEVICE,
OMP_LIST_CACHE,
- OMP_LIST_NUM
+ OMP_LIST_NUM,
+ OMP_LIST_LAST = OMP_LIST_NUM
};
/* Because a symbol can belong to multiple namelists, they must be
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 91e00c4..824ef79 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2713,6 +2713,29 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
return copy;
}
+/* Returns true if clause in list 'list' is compatible with any of
+ of the clauses in lists [0..list-1]. E.g., a reduction variable may
+ appear in both reduction and private clauses, so this function
+ will return true in this case. */
+
+static bool
+oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
+ gfc_symbol *sym, bool openacc)
+{
+ gfc_omp_namelist *n;
+
+ if (!openacc)
+ return false;
+
+ if (list != OMP_LIST_REDUCTION)
+ return false;
+
+ for (n = clauses->lists[OMP_LIST_FIRST]; n; n = n->next)
+ if (n->sym == sym)
+ return true;
+
+ return false;
+}
/* OpenMP directive resolving routines. */
@@ -2826,7 +2849,8 @@ resolve_omp_clauses (gfc_code *code, locus *where,
&& list != OMP_LIST_TO)
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark)
+ if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
+ n->sym, openacc))
gfc_error ("Symbol '%s' present on multiple clauses at %L",
n->sym->name, where);
else
@@ -3787,6 +3811,7 @@ struct omp_context
struct pointer_set_t *sharing_clauses;
struct pointer_set_t *private_iterators;
struct omp_context *previous;
+ bool is_openmp;
} *omp_current_ctx;
static gfc_code *omp_current_do_code;
static int omp_current_do_collapse;
@@ -3831,6 +3856,7 @@ gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns)
ctx.sharing_clauses = pointer_set_create ();
ctx.private_iterators = pointer_set_create ();
ctx.previous = omp_current_ctx;
+ ctx.is_openmp = true;
omp_current_ctx = &ctx;
for (list = 0; list < OMP_LIST_NUM; list++)
@@ -3925,7 +3951,12 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym)
if (omp_current_ctx == NULL)
return;
- if (pointer_set_contains (omp_current_ctx->sharing_clauses, sym))
+ /* An openacc context may represent a data clause. Abort if so. */
+ if (!omp_current_ctx->is_openmp && !oacc_is_loop (omp_current_ctx->code))
+ return;
+
+ if (omp_current_ctx->is_openmp
+ && pointer_set_contains (omp_current_ctx->sharing_clauses, sym))
return;
if (! pointer_set_insert (omp_current_ctx->private_iterators, sym))
@@ -4106,9 +4137,6 @@ resolve_omp_do (gfc_code *code)
}
}
-typedef struct omp_context oacc_context;
-oacc_context *oacc_current_ctx;
-
static bool
oacc_is_parallel (gfc_code *code)
{
@@ -4180,7 +4208,7 @@ switch (code->op)
static void
resolve_oacc_directive_inside_omp_region (gfc_code *code)
{
- if (omp_current_ctx != NULL)
+ if (omp_current_ctx != NULL && omp_current_ctx->is_openmp)
{
gfc_statement st = omp_code_to_statement (omp_current_ctx->code);
gfc_statement oacc_st = oacc_code_to_statement (code);
@@ -4193,9 +4221,9 @@ resolve_oacc_directive_inside_omp_region (gfc_code *code)
static void
resolve_omp_directive_inside_oacc_region (gfc_code *code)
{
- if (oacc_current_ctx != NULL)
+ if (omp_current_ctx != NULL && !omp_current_ctx->is_openmp)
{
- gfc_statement st = oacc_code_to_statement (oacc_current_ctx->code);
+ gfc_statement st = oacc_code_to_statement (omp_current_ctx->code);
gfc_statement omp_st = omp_code_to_statement (code);
gfc_error ("The %s directive cannot be specified within "
"a %s region at %L", gfc_ascii_statement (omp_st),
@@ -4282,12 +4310,12 @@ resolve_oacc_nested_loops (gfc_code *code, gfc_code* do_code, int collapse,
static void
resolve_oacc_params_in_parallel (gfc_code *code, const char *clause)
{
- oacc_context *c;
+ omp_context *c;
if (oacc_is_parallel (code))
gfc_error ("!$ACC LOOP %s in PARALLEL region doesn't allow "
"non-static arguments at %L", clause, &code->loc);
- for (c = oacc_current_ctx; c; c = c->previous)
+ for (c = omp_current_ctx; c; c = c->previous)
{
if (oacc_is_loop (c->code))
break;
@@ -4301,13 +4329,13 @@ resolve_oacc_params_in_parallel (gfc_code *code, const char *clause)
static void
resolve_oacc_loop_blocks (gfc_code *code)
{
- oacc_context *c;
+ omp_context *c;
if (!oacc_is_loop (code))
return;
if (code->op == EXEC_OACC_LOOP)
- for (c = oacc_current_ctx; c; c = c->previous)
+ for (c = omp_current_ctx; c; c = c->previous)
{
if (oacc_is_loop (c->code))
{
@@ -4419,17 +4447,21 @@ resolve_oacc_loop_blocks (gfc_code *code)
void
gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
{
- oacc_context ctx;
+ omp_context ctx;
resolve_oacc_loop_blocks (code);
ctx.code = code;
- ctx.previous = oacc_current_ctx;
- oacc_current_ctx = &ctx;
+ ctx.sharing_clauses = NULL;
+ ctx.private_iterators = pointer_set_create ();
+ ctx.previous = omp_current_ctx;
+ ctx.is_openmp = false;
+ omp_current_ctx = &ctx;
gfc_resolve_blocks (code->block, ns);
- oacc_current_ctx = ctx.previous;
+ pointer_set_destroy (ctx.private_iterators);
+ omp_current_ctx = ctx.previous;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
new file mode 100644
index 0000000..5aeee3b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
@@ -0,0 +1,38 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+end program test
+! { dg-prune-output "unimplemented" }
+! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-2.f95 b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
new file mode 100644
index 0000000..4b038f2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
@@ -0,0 +1,39 @@
+! { dg-do compile }
+
+! test for implicit private clauses in do loops
+
+program test
+ implicit none
+ integer :: i, j, k, a(10)
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop
+ do i = 1, 100
+ do j = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc data copy(a)
+
+ if(mod(1,10) .eq. 0) write(*,'(i5)') i
+
+ do i = 1, 100
+ !$acc parallel
+ !$acc loop
+ do j = 1, 100
+ do k = 1, 100
+ end do
+ end do
+ !$acc end parallel
+ end do
+
+ !$acc end data
+
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-3.f95 b/gcc/testsuite/gfortran.dg/goacc/private-3.f95
new file mode 100644
index 0000000..aa12a56
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-3.f95
@@ -0,0 +1,23 @@
+! { dg-do compile }
+
+! test for private variables in a reduction clause
+
+program test
+ implicit none
+ integer, parameter :: n = 100
+ integer :: i, k
+
+! FIXME: This causes an ICE in the gimplifier.
+! !$acc parallel private (k) reduction (+:k)
+! do i = 1, n
+! k = k + 1
+! end do
+! !$acc end parallel
+
+ !$acc parallel private (k)
+ !$acc loop reduction (+:k)
+ do i = 1, n
+ k = k + 1
+ end do
+ !$acc end parallel
+end program test
^ permalink raw reply [flat|nested] 12+ messages in thread
* [gomp4] don't error on implicitly private induction variables in gfortran
@ 2017-01-27 15:03 ` Cesar Philippidis
2017-01-27 15:07 ` Jakub Jelinek
2018-08-07 22:04 ` [PATCH][OpenACC] Don't " Cesar Philippidis
0 siblings, 2 replies; 12+ messages in thread
From: Cesar Philippidis @ 2017-01-27 15:03 UTC (permalink / raw)
To: gcc-patches, Fortran List
[-- Attachment #1: Type: text/plain, Size: 225 bytes --]
While experimenting with some new OpenACC benchmarks, I noticed that
gfortran errors when the user explicitly marks loop induction variables
as private. I applied this patch to gomp-4_0-branch to resolve that problem.
Cesar
[-- Attachment #2: gomp4-gfc-implifictly-private.diff --]
[-- Type: text/x-patch, Size: 1461 bytes --]
2017-01-26 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* openmp.c (gfc_resolve_oacc_blocks): Populate list of private
variables.
gcc/testsuite/
* gfortran.dg/goacc/implicitly-private.f90: New test.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 61940d7..2782a8d 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5192,7 +5192,8 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
{
fortran_omp_context ctx;
oacc_function dims = OACC_FUNCTION_NONE;
-
+ gfc_omp_namelist *n;
+
resolve_oacc_loop_blocks (code);
ctx.code = code;
@@ -5217,6 +5218,10 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
ctx.dims = dims;
omp_current_ctx = &ctx;
+ if (code->ext.omp_clauses)
+ for (n = code->ext.omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
+ ctx.private_iterators->add (n->sym);
+
gfc_resolve_blocks (code->block, ns);
omp_current_ctx = ctx.previous;
diff --git a/gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90 b/gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90
new file mode 100644
index 0000000..a687d8a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90
@@ -0,0 +1,12 @@
+! Ensure that implicitly private variables do not clash with those
+! that are explicitly private.
+
+program main
+ implicit none
+
+ integer i
+
+ !$acc parallel loop private(i)
+ do i = 1, 100
+ end do
+end program main
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [gomp4] don't error on implicitly private induction variables in gfortran
2017-01-27 15:03 ` [gomp4] don't error on implicitly private induction variables in gfortran Cesar Philippidis
@ 2017-01-27 15:07 ` Jakub Jelinek
2017-01-27 15:17 ` Cesar Philippidis
2018-08-07 22:04 ` [PATCH][OpenACC] Don't " Cesar Philippidis
1 sibling, 1 reply; 12+ messages in thread
From: Jakub Jelinek @ 2017-01-27 15:07 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches, Fortran List
On Fri, Jan 27, 2017 at 07:02:45AM -0800, Cesar Philippidis wrote:
> While experimenting with some new OpenACC benchmarks, I noticed that
> gfortran errors when the user explicitly marks loop induction variables
> as private. I applied this patch to gomp-4_0-branch to resolve that problem.
>
> Cesar
> 2017-01-26 Cesar Philippidis <cesar@codesourcery.com>
>
> gcc/fortran/
> * openmp.c (gfc_resolve_oacc_blocks): Populate list of private
> variables.
>
> gcc/testsuite/
> * gfortran.dg/goacc/implicitly-private.f90: New test.
>
>
> diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
> index 61940d7..2782a8d 100644
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -5192,7 +5192,8 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
> {
> fortran_omp_context ctx;
> oacc_function dims = OACC_FUNCTION_NONE;
> -
> + gfc_omp_namelist *n;
> +
> resolve_oacc_loop_blocks (code);
>
> ctx.code = code;
> @@ -5217,6 +5218,10 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
> ctx.dims = dims;
> omp_current_ctx = &ctx;
>
> + if (code->ext.omp_clauses)
> + for (n = code->ext.omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
> + ctx.private_iterators->add (n->sym);
Do you really want to add to private_iterators even variables that aren't
iterators, just arbitrary private variable?
Jakub
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [gomp4] don't error on implicitly private induction variables in gfortran
2017-01-27 15:07 ` Jakub Jelinek
@ 2017-01-27 15:17 ` Cesar Philippidis
0 siblings, 0 replies; 12+ messages in thread
From: Cesar Philippidis @ 2017-01-27 15:17 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches, Fortran List
On 01/27/2017 07:07 AM, Jakub Jelinek wrote:
> On Fri, Jan 27, 2017 at 07:02:45AM -0800, Cesar Philippidis wrote:
>> diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
>> index 61940d7..2782a8d 100644
>> --- a/gcc/fortran/openmp.c
>> +++ b/gcc/fortran/openmp.c
>> @@ -5192,7 +5192,8 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
>> {
>> fortran_omp_context ctx;
>> oacc_function dims = OACC_FUNCTION_NONE;
>> -
>> + gfc_omp_namelist *n;
>> +
>> resolve_oacc_loop_blocks (code);
>>
>> ctx.code = code;
>> @@ -5217,6 +5218,10 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
>> ctx.dims = dims;
>> omp_current_ctx = &ctx;
>>
>> + if (code->ext.omp_clauses)
>> + for (n = code->ext.omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
>> + ctx.private_iterators->add (n->sym);
>
> Do you really want to add to private_iterators even variables that aren't
> iterators, just arbitrary private variable?
Well, the other thing we can do is scan the list of private clauses and
look for the induction variable. But that would result in an O(n) lookup
vs O(1) with the hash. My impression is that this approach should be
innocuous, because it doesn't appear that private_iterators are used for
anything else at the moment.
Cesar
^ permalink raw reply [flat|nested] 12+ messages in thread
* [PATCH][OpenACC] Don't error on implicitly private induction variables in gfortran
@ 2018-08-07 22:04 ` Cesar Philippidis
2019-04-17 8:38 ` [PR90048] Fortran OpenACC 'private' clause rejected for predetermined private loop iteration variable (was: [patch,gomp4] make fortran loop variables implicitly private in openacc) Thomas Schwinge
0 siblings, 1 reply; 12+ messages in thread
From: Cesar Philippidis @ 2018-08-07 22:04 UTC (permalink / raw)
To: Fortran List, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 382 bytes --]
At present, the fortran FE reports an error if the user adds an explicit
private clause to an induction variable used by an acc loop. This patch
teaches the fortran acc block resolver how to cope with "duplicate"
private clauses, so that it doesn't error anymore.
Is this patch OK for trunk? I bootstrapped and regression tested it for
x86_64 with nvptx offloading.
Thanks,
Cesar
[-- Attachment #2: 0001-OpenACC-Don-t-error-on-implicitly-private-induction-.patch --]
[-- Type: text/x-patch, Size: 1924 bytes --]
From 576b2a7d5574400f067ec309929b38b324d8c6f6 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Fri, 27 Jan 2017 14:58:16 +0000
Subject: [PATCH] [OpenACC] Don't error on implicitly private induction
variables in gfortran
2018-XX-YY Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* openmp.c (gfc_resolve_oacc_blocks): Populate list of private
variables.
gcc/testsuite/
* gfortran.dg/goacc/implicitly-private.f90: New test.
---
gcc/fortran/openmp.c | 5 +++++
gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90 | 12 ++++++++++++
2 files changed, 17 insertions(+)
create mode 100644 gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index b346b51..798c5fa 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5951,6 +5951,7 @@ void
gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
{
fortran_omp_context ctx;
+ gfc_omp_namelist *n;
resolve_oacc_loop_blocks (code);
@@ -5961,6 +5962,10 @@ gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
ctx.is_openmp = false;
omp_current_ctx = &ctx;
+ if (code->ext.omp_clauses)
+ for (n = code->ext.omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next)
+ ctx.private_iterators->add (n->sym);
+
gfc_resolve_blocks (code->block, ns);
omp_current_ctx = ctx.previous;
diff --git a/gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90 b/gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90
new file mode 100644
index 0000000..a687d8a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/implicitly-private.f90
@@ -0,0 +1,12 @@
+! Ensure that implicitly private variables do not clash with those
+! that are explicitly private.
+
+program main
+ implicit none
+
+ integer i
+
+ !$acc parallel loop private(i)
+ do i = 1, 100
+ end do
+end program main
--
2.7.4
^ permalink raw reply [flat|nested] 12+ messages in thread
* [PR90048] Fortran OpenACC 'private' clause rejected for predetermined private loop iteration variable (was: [patch,gomp4] make fortran loop variables implicitly private in openacc)
2018-08-07 22:04 ` [PATCH][OpenACC] Don't " Cesar Philippidis
@ 2019-04-17 8:38 ` Thomas Schwinge
0 siblings, 0 replies; 12+ messages in thread
From: Thomas Schwinge @ 2019-04-17 8:38 UTC (permalink / raw)
To: fortran, gcc-patches; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 2475 bytes --]
Hi!
On Mon, 11 Aug 2014 16:55:28 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> According to section 2.6.1 in the openacc spec, fortran loop variables
> should be implicitly private like in openmp.
More correctly, they are "predetermined private" (which cannot be
overridden), not "implicit private" (which could be overridden with a
different explicit clause).
> This patch does just so.
But it also introduced PR90048 "Fortran OpenACC 'private' clause rejected
for predetermined private loop iteration variable".
Instead of the patch "don't error on implicitly private induction
variables in gfortran" proposed by Cesar, and then challenged by Jakub, I
have now committed a different patch (more similar to the existing
handling for OpenMP) to trunk in r270406 "[PR90048] Fortran OpenACC
'private' clause rejected for predetermined private loop iteration
variable", see attached.
I have a cleanup patch (for next GCC development stage 1), which will
simply merge the special-case 'gfc_resolve_oacc_blocks' into the generic
'gfc_resolve_omp_parallel_blocks', see attached.
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
> @@ -0,0 +1,39 @@
> +! { dg-do compile }
> +! { dg-additional-options "-fdump-tree-omplower" }
> +
> +! test for implicit private clauses in do loops
> +
> +program test
> + implicit none
> + integer :: i, j, k
> + logical :: l
> +
> + !$acc parallel
> + !$acc loop
> + do i = 1, 100
> + end do
> + !$acc end parallel
> +
> + !$acc parallel
> + !$acc loop
> + do i = 1, 100
> + do j = 1, 100
> + end do
> + end do
> + !$acc end parallel
> +
> + !$acc parallel
> + !$acc loop
> + do i = 1, 100
> + do j = 1, 100
> + do k = 1, 100
> + end do
> + end do
> + end do
> + !$acc end parallel
> +end program test
> +! { dg-prune-output "unimplemented" }
> +! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } }
> +! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
> +! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
> +! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
I turned that one and 'gfortran.dg/goacc/private-2.f95' into more
elaborate testcases, committed to trunk in r270405 "[PR90067, PR90114]
Document Fortran OpenACC predetermined private status quo", see attached.
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-PR90048-Fortran-OpenACC-private-clause-rejecte.trunk.patch --]
[-- Type: text/x-diff, Size: 36126 bytes --]
From b8d03885017763f914a48b19b6cb383239430b97 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 17 Apr 2019 08:34:20 +0000
Subject: [PATCH] [PR90048] Fortran OpenACC 'private' clause rejected for
predetermined private loop iteration variable
gcc/fortran/
PR fortran/90048
* openmp.c (gfc_resolve_do_iterator): Handle sharing_clauses for
OpenACC, too.
(gfc_resolve_oacc_blocks): Populate sharing_clauses with private
clauses.
gcc/testsuite/
PR fortran/90048
* gfortran.dg/goacc/private-explicit-kernels-1.f95: New file.
* gfortran.dg/goacc/private-explicit-parallel-1.f95: Likewise.
* gfortran.dg/goacc/private-explicit-routine-1.f95: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@270406 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/fortran/ChangeLog | 8 +
gcc/fortran/openmp.c | 20 +-
gcc/testsuite/ChangeLog | 5 +
.../goacc/private-explicit-kernels-1.f95 | 248 ++++++++++++++++++
.../goacc/private-explicit-parallel-1.f95 | 247 +++++++++++++++++
.../goacc/private-explicit-routine-1.f95 | 146 +++++++++++
6 files changed, 671 insertions(+), 3 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-explicit-parallel-1.f95
create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-explicit-routine-1.f95
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index e27743cac280..1ff03e1e85b5 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,11 @@
+2019-04-17 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR fortran/90048
+ * openmp.c (gfc_resolve_do_iterator): Handle sharing_clauses for
+ OpenACC, too.
+ (gfc_resolve_oacc_blocks): Populate sharing_clauses with private
+ clauses.
+
2019-04-14 Paul Thomas <pault@gcc.gnu.org>
PR fortran/89843
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 9fc236760a1c..1c7bce6c3000 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5510,8 +5510,7 @@ gfc_resolve_do_iterator (gfc_code *code, gfc_symbol *sym, bool add_clause)
if (!omp_current_ctx->is_openmp && !oacc_is_loop (omp_current_ctx->code))
return;
- if (omp_current_ctx->is_openmp
- && omp_current_ctx->sharing_clauses->contains (sym))
+ if (omp_current_ctx->sharing_clauses->contains (sym))
return;
if (! omp_current_ctx->private_iterators->add (sym) && add_clause)
@@ -5971,19 +5970,34 @@ void
gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
{
fortran_omp_context ctx;
+ gfc_omp_clauses *omp_clauses = code->ext.omp_clauses;
+ gfc_omp_namelist *n;
+ int list;
resolve_oacc_loop_blocks (code);
ctx.code = code;
- ctx.sharing_clauses = NULL;
+ ctx.sharing_clauses = new hash_set<gfc_symbol *>;
ctx.private_iterators = new hash_set<gfc_symbol *>;
ctx.previous = omp_current_ctx;
ctx.is_openmp = false;
omp_current_ctx = &ctx;
+ for (list = 0; list < OMP_LIST_NUM; list++)
+ switch (list)
+ {
+ case OMP_LIST_PRIVATE:
+ for (n = omp_clauses->lists[list]; n; n = n->next)
+ ctx.sharing_clauses->add (n->sym);
+ break;
+ default:
+ break;
+ }
+
gfc_resolve_blocks (code->block, ns);
omp_current_ctx = ctx.previous;
+ delete ctx.sharing_clauses;
delete ctx.private_iterators;
}
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 0ed8b3c52837..937057421e45 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,5 +1,10 @@
2019-04-17 Thomas Schwinge <thomas@codesourcery.com>
+ PR fortran/90048
+ * gfortran.dg/goacc/private-explicit-kernels-1.f95: New file.
+ * gfortran.dg/goacc/private-explicit-parallel-1.f95: Likewise.
+ * gfortran.dg/goacc/private-explicit-routine-1.f95: Likewise.
+
PR fortran/90067
PR fortran/90114
* gfortran.dg/goacc/private-1.f95: Remove file.
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
new file mode 100644
index 000000000000..5d563d226b0c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
@@ -0,0 +1,248 @@
+! Explicit 'private' clauses related to 'do' loops inside an OpenACC
+! 'kernels' construct.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+program test
+ implicit none
+ integer :: i0_1
+ integer :: i0_2, j0_2
+ integer :: i1_s
+ integer :: i1_c
+ integer :: i2_1_s, j2_1_s
+ integer :: i2_1_c, j2_1_c
+ integer :: i2_2_s, j2_2_s
+ integer :: i2_3_s, j2_3_s
+ integer :: i2_3_c, j2_3_c
+ integer :: i3_1_s, j3_1_s, k3_1_s
+ integer :: i3_1_c, j3_1_c, k3_1_c
+ integer :: i3_2_s, j3_2_s, k3_2_s
+ integer :: i3_2_c, j3_2_c, k3_2_c
+ integer :: i3_3_s, j3_3_s, k3_3_s
+ integer :: i3_3_c, j3_3_c, k3_3_c
+ integer :: i3_4_s, j3_4_s, k3_4_s
+ integer :: i3_4_c, j3_4_c, k3_4_c
+ integer :: i3_5_s, j3_5_s, k3_5_s
+
+ !$acc kernels ! Explicit "private(i0_1)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i0_1 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i0_1 = 1, 100
+ end do
+ !$acc end kernels
+
+ !$acc kernels ! Explicit "private(i0_2, j0_2)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:j0_2 \\\[len: \[0-9\]+\\\]\\) map\\(force_tofrom:i0_2 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i0_2 = 1, 100
+ do j0_2 = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels
+ !$acc loop private(i1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+ do i1_s = 1, 100
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i1_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "gimple" } }
+ do i1_c = 1, 100
+ end do
+
+ !$acc kernels
+ !$acc loop private(i2_1_s, j2_1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+ do i2_1_s = 1, 100
+ do j2_1_s = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i2_1_c, j2_1_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+ do i2_1_c = 1, 100
+ do j2_1_c = 1, 100
+ end do
+ end do
+
+ !$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i2_2_s = 1, 100
+ !$acc loop private(j2_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+ do j2_2_s = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels
+ !$acc loop private(i2_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+ do i2_3_s = 1, 100
+ !$acc loop private(j2_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+ do j2_3_s = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i2_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "gimple" } }
+ do i2_3_c = 1, 100
+ !$acc loop private(j2_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "gimple" } }
+ do j2_3_c = 1, 100
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop private(i3_1_s, j3_1_s, k3_1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+ do i3_1_s = 1, 100
+ do j3_1_s = 1, 100
+ do k3_1_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i3_1_c, j3_1_c, k3_1_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+ do i3_1_c = 1, 100
+ do j3_1_c = 1, 100
+ do k3_1_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop private(i3_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+ do i3_2_s = 1, 100
+ !$acc loop private(j3_2_s, k3_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+ do j3_2_s = 1, 100
+ do k3_2_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i3_2_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "gimple" } }
+ do i3_2_c = 1, 100
+ !$acc loop private(j3_2_c, k3_2_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "gimple" } }
+ do j3_2_c = 1, 100
+ do k3_2_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop private(i3_3_s, j3_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+ do i3_3_s = 1, 100
+ do j3_3_s = 1, 100
+ !$acc loop private(k3_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+ do k3_3_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i3_3_c, j3_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "gimple" } }
+ do i3_3_c = 1, 100
+ do j3_3_c = 1, 100
+ !$acc loop private(k3_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "gimple" } }
+ do k3_3_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop private(i3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+ do i3_4_s = 1, 100
+ !$acc loop private(j3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+ do j3_4_s = 1, 100
+ !$acc loop private(k3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+ do k3_4_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop private(i3_4_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "gimple" } }
+ do i3_4_c = 1, 100
+ !$acc loop private(j3_4_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "gimple" } }
+ do j3_4_c = 1, 100
+ !$acc loop private(k3_4_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "gimple" } }
+ do k3_4_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i3_5_s = 1, 100
+ !$acc loop private(j3_5_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+ do j3_5_s = 1, 100
+ !$acc loop private(k3_5_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+ do k3_5_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-parallel-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-parallel-1.f95
new file mode 100644
index 000000000000..bac919b645de
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-parallel-1.f95
@@ -0,0 +1,247 @@
+! Explicit 'private' clauses related to 'do' loops inside an OpenACC
+! 'parallel' construct.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+program test
+ implicit none
+ integer :: i0_1
+ integer :: i0_2, j0_2
+ integer :: i1_s
+ integer :: i1_c
+ integer :: i2_1_s, j2_1_s
+ integer :: i2_1_c, j2_1_c
+ integer :: i2_2_s, j2_2_s
+ integer :: i2_3_s, j2_3_s
+ integer :: i2_3_c, j2_3_c
+ integer :: i3_1_s, j3_1_s, k3_1_s
+ integer :: i3_1_c, j3_1_c, k3_1_c
+ integer :: i3_2_s, j3_2_s, k3_2_s
+ integer :: i3_2_c, j3_2_c, k3_2_c
+ integer :: i3_3_s, j3_3_s, k3_3_s
+ integer :: i3_3_c, j3_3_c, k3_3_c
+ integer :: i3_4_s, j3_4_s, k3_4_s
+ integer :: i3_4_c, j3_4_c, k3_4_c
+ integer :: i3_5_s, j3_5_s, k3_5_s
+
+ !$acc parallel private(i0_1)
+ ! { dg-final { scan-tree-dump-times "#pragma acc parallel private\\(i0_1\\)" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel private\\(i0_1\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i0_1\\)" 0 "gimple" } }
+ do i0_1 = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel private(i0_2, j0_2)
+ ! { dg-final { scan-tree-dump-times "#pragma acc parallel private\\(i0_2\\) private\\(j0_2\\)" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel private\\(i0_2\\) private\\(j0_2\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i0_2\\)" 0 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(j0_2\\)" 0 "gimple" } }
+ do i0_2 = 1, 100
+ do j0_2 = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop private(i1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+ do i1_s = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i1_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "gimple" } }
+ do i1_c = 1, 100
+ end do
+
+ !$acc parallel
+ !$acc loop private(i2_1_s, j2_1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+ do i2_1_s = 1, 100
+ do j2_1_s = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i2_1_c, j2_1_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+ do i2_1_c = 1, 100
+ do j2_1_c = 1, 100
+ end do
+ end do
+
+ !$acc parallel private(i2_2_s)
+ ! { dg-final { scan-tree-dump-times "#pragma acc parallel private\\(i2_2_s\\)" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel private\\(i2_2_s\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i2_2_s\\)" 0 "gimple" } }
+ do i2_2_s = 1, 100
+ !$acc loop private(j2_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+ do j2_2_s = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop private(i2_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+ do i2_3_s = 1, 100
+ !$acc loop private(j2_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+ do j2_3_s = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i2_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "gimple" } }
+ do i2_3_c = 1, 100
+ !$acc loop private(j2_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "gimple" } }
+ do j2_3_c = 1, 100
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop private(i3_1_s, j3_1_s, k3_1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+ do i3_1_s = 1, 100
+ do j3_1_s = 1, 100
+ do k3_1_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i3_1_c, j3_1_c, k3_1_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+ do i3_1_c = 1, 100
+ do j3_1_c = 1, 100
+ do k3_1_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop private(i3_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+ do i3_2_s = 1, 100
+ !$acc loop private(j3_2_s, k3_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+ do j3_2_s = 1, 100
+ do k3_2_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i3_2_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "gimple" } }
+ do i3_2_c = 1, 100
+ !$acc loop private(j3_2_c, k3_2_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "gimple" } }
+ do j3_2_c = 1, 100
+ do k3_2_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop private(i3_3_s, j3_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+ do i3_3_s = 1, 100
+ do j3_3_s = 1, 100
+ !$acc loop private(k3_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+ do k3_3_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i3_3_c, j3_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "gimple" } }
+ do i3_3_c = 1, 100
+ do j3_3_c = 1, 100
+ !$acc loop private(k3_3_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "gimple" } }
+ do k3_3_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop private(i3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+ do i3_4_s = 1, 100
+ !$acc loop private(j3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+ do j3_4_s = 1, 100
+ !$acc loop private(k3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+ do k3_4_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop private(i3_4_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "gimple" } }
+ do i3_4_c = 1, 100
+ !$acc loop private(j3_4_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "gimple" } }
+ do j3_4_c = 1, 100
+ !$acc loop private(k3_4_c) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "gimple" } }
+ do k3_4_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel private(i3_5_s)
+ ! { dg-final { scan-tree-dump-times "#pragma acc parallel private\\(i3_5_s\\)" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel private\\(i3_5_s\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i3_5_s\\)" 0 "gimple" } }
+ do i3_5_s = 1, 100
+ !$acc loop private(j3_5_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+ do j3_5_s = 1, 100
+ !$acc loop private(k3_5_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+ do k3_5_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-routine-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-routine-1.f95
new file mode 100644
index 000000000000..b7f8c3d5b378
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-routine-1.f95
@@ -0,0 +1,146 @@
+! Explicit 'private' clauses related to 'do' loops inside an OpenACC
+! accelerator routine.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+! The PR90114 XFAILs need to scan for the appropriate predetermined private
+! level.
+
+subroutine test
+ implicit none
+ integer :: i0_1
+ integer :: i0_2, j0_2
+ integer :: i1_s
+ integer :: i2_1_s, j2_1_s
+ integer :: i2_2_s, j2_2_s
+ integer :: i2_3_s, j2_3_s
+ integer :: i3_1_s, j3_1_s, k3_1_s
+ integer :: i3_2_s, j3_2_s, k3_2_s
+ integer :: i3_3_s, j3_3_s, k3_3_s
+ integer :: i3_4_s, j3_4_s, k3_4_s
+ integer :: i3_5_s, j3_5_s, k3_5_s
+ !$acc routine gang
+
+ ! Explicit "private(i0_1)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i0_1 = 1, 100
+ end do
+
+ ! Explicit "private(i0_2, j0_2)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i0_2 = 1, 100
+ do j0_2 = 1, 100
+ end do
+ end do
+
+ !$acc loop private(i1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+ do i1_s = 1, 100
+ end do
+
+ !$acc loop private(i2_1_s, j2_1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+ do i2_1_s = 1, 100
+ do j2_1_s = 1, 100
+ end do
+ end do
+
+ ! Explicit "private(i2_2_s)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i2_2_s = 1, 100
+ !$acc loop private(j2_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+ do j2_2_s = 1, 100
+ end do
+ end do
+
+ !$acc loop private(i2_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+ do i2_3_s = 1, 100
+ !$acc loop private(j2_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+ do j2_3_s = 1, 100
+ end do
+ end do
+
+ !$acc loop private(i3_1_s, j3_1_s, k3_1_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+ do i3_1_s = 1, 100
+ do j3_1_s = 1, 100
+ do k3_1_s = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc loop private(i3_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+ do i3_2_s = 1, 100
+ !$acc loop private(j3_2_s, k3_2_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+ do j3_2_s = 1, 100
+ do k3_2_s = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc loop private(i3_3_s, j3_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+ do i3_3_s = 1, 100
+ do j3_3_s = 1, 100
+ !$acc loop private(k3_3_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+ do k3_3_s = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc loop private(i3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+ do i3_4_s = 1, 100
+ !$acc loop private(j3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+ do j3_4_s = 1, 100
+ !$acc loop private(k3_4_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+ do k3_4_s = 1, 100
+ end do
+ end do
+ end do
+
+ ! Explicit "private(i3_5_s)" clause cannot be specified here.
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i3_5_s = 1, 100
+ !$acc loop private(j3_5_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+ do j3_5_s = 1, 100
+ !$acc loop private(k3_5_s) independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+ do k3_5_s = 1, 100
+ end do
+ end do
+ end do
+end subroutine test
--
2.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-Fortran-Clean-up-gfc_resolve_oacc_blocks.patch --]
[-- Type: text/x-diff, Size: 5825 bytes --]
From e2a781a0e6519ca5b135e697668dbc85dea1fe0f Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 17 Apr 2019 09:22:06 +0200
Subject: [PATCH] [Fortran] Clean up 'gfc_resolve_oacc_blocks'
Developed in context of PR90048.
---
gcc/fortran/gfortran.h | 12 ++++++++--
gcc/fortran/openmp.c | 52 +++++++++++-------------------------------
gcc/fortran/resolve.c | 6 ++---
3 files changed, 26 insertions(+), 44 deletions(-)
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index d4426908ed2a..3dc8d57c78b5 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -671,6 +671,14 @@ enum gfc_param_spec_type
SPEC_DEFERRED
};
+/* Kind of an OMP construct. */
+
+enum omp_kind
+{
+ OMP_OPENACC,
+ OMP_OPENMP
+};
+
/************************* Structures *****************************/
/* Used for keeping things in balanced binary trees. */
@@ -3196,7 +3204,8 @@ gfc_omp_udr *gfc_omp_udr_find (gfc_symtree *, gfc_typespec *);
void gfc_resolve_omp_directive (gfc_code *, gfc_namespace *);
void gfc_resolve_do_iterator (gfc_code *, gfc_symbol *, bool);
void gfc_resolve_omp_local_vars (gfc_namespace *);
-void gfc_resolve_omp_parallel_blocks (gfc_code *, gfc_namespace *);
+void gfc_resolve_omp_parallel_blocks (gfc_code *, gfc_namespace *,
+ enum omp_kind);
void gfc_resolve_omp_do_blocks (gfc_code *, gfc_namespace *);
void gfc_resolve_omp_declare_simd (gfc_namespace *);
void gfc_resolve_omp_udrs (gfc_symtree *);
@@ -3206,7 +3215,6 @@ void gfc_free_expr_list (gfc_expr_list *);
void gfc_resolve_oacc_directive (gfc_code *, gfc_namespace *);
void gfc_resolve_oacc_declare (gfc_namespace *);
void gfc_resolve_oacc_parallel_loop_blocks (gfc_code *, gfc_namespace *);
-void gfc_resolve_oacc_blocks (gfc_code *, gfc_namespace *);
void gfc_resolve_oacc_routines (gfc_namespace *);
/* expr.c */
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index e372326db1aa..414dd5e355e5 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5432,9 +5432,15 @@ gfc_resolve_omp_do_blocks (gfc_code *code, gfc_namespace *ns)
}
+static void resolve_oacc_loop_blocks (gfc_code *);
+
void
-gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns)
+gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns,
+ enum omp_kind omp_kind)
{
+ if (omp_kind == OMP_OPENACC)
+ resolve_oacc_loop_blocks (code);
+
struct fortran_omp_context ctx;
gfc_omp_clauses *omp_clauses = code->ext.omp_clauses;
gfc_omp_namelist *n;
@@ -5444,18 +5450,21 @@ gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns)
ctx.sharing_clauses = new hash_set<gfc_symbol *>;
ctx.private_iterators = new hash_set<gfc_symbol *>;
ctx.previous = omp_current_ctx;
- ctx.is_openmp = true;
+ ctx.is_openmp = (omp_kind == OMP_OPENMP);
omp_current_ctx = &ctx;
for (list = 0; list < OMP_LIST_NUM; list++)
switch (list)
{
case OMP_LIST_SHARED:
- case OMP_LIST_PRIVATE:
case OMP_LIST_FIRSTPRIVATE:
case OMP_LIST_LASTPRIVATE:
case OMP_LIST_REDUCTION:
case OMP_LIST_LINEAR:
+ if (omp_kind == OMP_OPENACC)
+ continue;
+ /* FALLTHRU */
+ case OMP_LIST_PRIVATE:
for (n = omp_clauses->lists[list]; n; n = n->next)
ctx.sharing_clauses->add (n->sym);
break;
@@ -5479,6 +5488,7 @@ gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns)
case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ gcc_checking_assert (omp_kind == OMP_OPENMP);
gfc_resolve_omp_do_blocks (code, ns);
break;
default:
@@ -5976,42 +5986,6 @@ resolve_oacc_loop_blocks (gfc_code *code)
}
-void
-gfc_resolve_oacc_blocks (gfc_code *code, gfc_namespace *ns)
-{
- fortran_omp_context ctx;
- gfc_omp_clauses *omp_clauses = code->ext.omp_clauses;
- gfc_omp_namelist *n;
- int list;
-
- resolve_oacc_loop_blocks (code);
-
- ctx.code = code;
- ctx.sharing_clauses = new hash_set<gfc_symbol *>;
- ctx.private_iterators = new hash_set<gfc_symbol *>;
- ctx.previous = omp_current_ctx;
- ctx.is_openmp = false;
- omp_current_ctx = &ctx;
-
- for (list = 0; list < OMP_LIST_NUM; list++)
- switch (list)
- {
- case OMP_LIST_PRIVATE:
- for (n = omp_clauses->lists[list]; n; n = n->next)
- ctx.sharing_clauses->add (n->sym);
- break;
- default:
- break;
- }
-
- gfc_resolve_blocks (code->block, ns);
-
- omp_current_ctx = ctx.previous;
- delete ctx.sharing_clauses;
- delete ctx.private_iterators;
-}
-
-
static void
resolve_oacc_loop (gfc_code *code)
{
diff --git a/gcc/fortran/resolve.c b/gcc/fortran/resolve.c
index cb41da08526f..8ee3f9fb3afb 100644
--- a/gcc/fortran/resolve.c
+++ b/gcc/fortran/resolve.c
@@ -11291,12 +11291,12 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns)
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
- gfc_resolve_oacc_blocks (code, ns);
+ gfc_resolve_omp_parallel_blocks (code, ns, OMP_OPENACC);
break;
case EXEC_OMP_PARALLEL_WORKSHARE:
omp_workshare_save = omp_workshare_flag;
omp_workshare_flag = 1;
- gfc_resolve_omp_parallel_blocks (code, ns);
+ gfc_resolve_omp_parallel_blocks (code, ns, OMP_OPENMP);
break;
case EXEC_OMP_PARALLEL:
case EXEC_OMP_PARALLEL_DO:
@@ -11320,7 +11320,7 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns)
case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
omp_workshare_save = omp_workshare_flag;
omp_workshare_flag = 0;
- gfc_resolve_omp_parallel_blocks (code, ns);
+ gfc_resolve_omp_parallel_blocks (code, ns, OMP_OPENMP);
break;
case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DISTRIBUTE_SIMD:
--
2.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #4: 0001-PR90067-PR90114-Document-Fortran-OpenACC-prede.trunk.patch --]
[-- Type: text/x-diff, Size: 35230 bytes --]
From 27d00f9196a9ca0d48b0f0934b715307c4b08f5d Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 17 Apr 2019 08:34:10 +0000
Subject: [PATCH] [PR90067, PR90114] Document Fortran OpenACC predetermined
private status quo
gcc/testsuite/
PR fortran/90067
PR fortran/90114
* gfortran.dg/goacc/private-1.f95: Remove file.
* gfortran.dg/goacc/private-2.f95: Likewise.
* gfortran.dg/goacc/private-predetermined-kernels-1.f95: New file.
* gfortran.dg/goacc/private-predetermined-parallel-1.f95:
Likewise.
* gfortran.dg/goacc/private-predetermined-routine-1.f95: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@270405 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/testsuite/ChangeLog | 11 +
gcc/testsuite/gfortran.dg/goacc/private-1.f95 | 37 ---
gcc/testsuite/gfortran.dg/goacc/private-2.f95 | 39 ---
.../goacc/private-predetermined-kernels-1.f95 | 248 +++++++++++++++++
.../private-predetermined-parallel-1.f95 | 253 ++++++++++++++++++
.../goacc/private-predetermined-routine-1.f95 | 142 ++++++++++
6 files changed, 654 insertions(+), 76 deletions(-)
delete mode 100644 gcc/testsuite/gfortran.dg/goacc/private-1.f95
delete mode 100644 gcc/testsuite/gfortran.dg/goacc/private-2.f95
create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95
create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 23f77ec781a6..0ed8b3c52837 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,14 @@
+2019-04-17 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR fortran/90067
+ PR fortran/90114
+ * gfortran.dg/goacc/private-1.f95: Remove file.
+ * gfortran.dg/goacc/private-2.f95: Likewise.
+ * gfortran.dg/goacc/private-predetermined-kernels-1.f95: New file.
+ * gfortran.dg/goacc/private-predetermined-parallel-1.f95:
+ Likewise.
+ * gfortran.dg/goacc/private-predetermined-routine-1.f95: Likewise.
+
2019-04-17 Jakub Jelinek <jakub@redhat.com>
PR target/89093
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
deleted file mode 100644
index 23ce95ad8d24..000000000000
--- a/gcc/testsuite/gfortran.dg/goacc/private-1.f95
+++ /dev/null
@@ -1,37 +0,0 @@
-! { dg-do compile }
-! { dg-additional-options "-fdump-tree-omplower" }
-
-! test for implicit private clauses in do loops
-
-program test
- implicit none
- integer :: i, j, k
-
- !$acc parallel
- !$acc loop
- do i = 1, 100
- end do
- !$acc end parallel
-
- !$acc parallel
- !$acc loop
- do i = 1, 100
- do j = 1, 100
- end do
- end do
- !$acc end parallel
-
- !$acc parallel
- !$acc loop
- do i = 1, 100
- do j = 1, 100
- do k = 1, 100
- end do
- end do
- end do
- !$acc end parallel
-end program test
-! { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel" 3 "omplower" } }
-! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
-! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-2.f95 b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
deleted file mode 100644
index 4b038f2b5f20..000000000000
--- a/gcc/testsuite/gfortran.dg/goacc/private-2.f95
+++ /dev/null
@@ -1,39 +0,0 @@
-! { dg-do compile }
-
-! test for implicit private clauses in do loops
-
-program test
- implicit none
- integer :: i, j, k, a(10)
-
- !$acc parallel
- !$acc loop
- do i = 1, 100
- end do
- !$acc end parallel
-
- !$acc parallel
- !$acc loop
- do i = 1, 100
- do j = 1, 100
- end do
- end do
- !$acc end parallel
-
- !$acc data copy(a)
-
- if(mod(1,10) .eq. 0) write(*,'(i5)') i
-
- do i = 1, 100
- !$acc parallel
- !$acc loop
- do j = 1, 100
- do k = 1, 100
- end do
- end do
- !$acc end parallel
- end do
-
- !$acc end data
-
-end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
new file mode 100644
index 000000000000..12a7854526a9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -0,0 +1,248 @@
+! Predetermined 'private' clauses related to 'do' loops inside an OpenACC
+! 'kernels' construct.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+program test
+ implicit none
+ integer :: i0_1
+ integer :: i0_2, j0_2
+ integer :: i1_s
+ integer :: i1_c
+ integer :: i2_1_s, j2_1_s
+ integer :: i2_1_c, j2_1_c
+ integer :: i2_2_s, j2_2_s
+ integer :: i2_3_s, j2_3_s
+ integer :: i2_3_c, j2_3_c
+ integer :: i3_1_s, j3_1_s, k3_1_s
+ integer :: i3_1_c, j3_1_c, k3_1_c
+ integer :: i3_2_s, j3_2_s, k3_2_s
+ integer :: i3_2_c, j3_2_c, k3_2_c
+ integer :: i3_3_s, j3_3_s, k3_3_s
+ integer :: i3_3_c, j3_3_c, k3_3_c
+ integer :: i3_4_s, j3_4_s, k3_4_s
+ integer :: i3_4_c, j3_4_c, k3_4_c
+ integer :: i3_5_s, j3_5_s, k3_5_s
+
+ !$acc kernels
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i0_1 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i0_1 = 1, 100
+ end do
+ !$acc end kernels
+
+ !$acc kernels
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:j0_2 \\\[len: \[0-9\]+\\\]\\) map\\(force_tofrom:i0_2 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i0_2 = 1, 100
+ do j0_2 = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+ do i1_s = 1, 100
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "gimple" } }
+ do i1_c = 1, 100
+ end do
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+ do i2_1_s = 1, 100
+ do j2_1_s = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+ do i2_1_c = 1, 100
+ do j2_1_c = 1, 100
+ end do
+ end do
+
+ !$acc kernels
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i2_2_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+ do j2_2_s = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+ do i2_3_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+ do j2_3_s = 1, 100
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "gimple" } }
+ do i2_3_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "gimple" } }
+ do j2_3_c = 1, 100
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+ do i3_1_s = 1, 100
+ do j3_1_s = 1, 100
+ do k3_1_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+ do i3_1_c = 1, 100
+ do j3_1_c = 1, 100
+ do k3_1_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+ do i3_2_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+ do j3_2_s = 1, 100
+ do k3_2_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "gimple" } }
+ do i3_2_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "gimple" } }
+ do j3_2_c = 1, 100
+ do k3_2_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+ do i3_3_s = 1, 100
+ do j3_3_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+ do k3_3_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "gimple" } }
+ do i3_3_c = 1, 100
+ do j3_3_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "gimple" } }
+ do k3_3_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+ do i3_4_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+ do j3_4_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+ do k3_4_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "gimple" } }
+ do i3_4_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "gimple" } }
+ do j3_4_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "gimple" } }
+ do k3_4_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc kernels
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i3_5_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+ do j3_5_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+ do k3_5_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end kernels
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95
new file mode 100644
index 000000000000..a2fa71a58b72
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95
@@ -0,0 +1,253 @@
+! Predetermined 'private' clauses related to 'do' loops inside an OpenACC
+! 'parallel' construct.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+program test
+ implicit none
+ integer :: i0_1
+ integer :: i0_2, j0_2
+ integer :: i1_s
+ integer :: i1_c
+ integer :: i2_1_s, j2_1_s
+ integer :: i2_1_c, j2_1_c
+ integer :: i2_2_s, j2_2_s
+ integer :: i2_3_s, j2_3_s
+ integer :: i2_3_c, j2_3_c
+ integer :: i3_1_s, j3_1_s, k3_1_s
+ integer :: i3_1_c, j3_1_c, k3_1_c
+ integer :: i3_2_s, j3_2_s, k3_2_s
+ integer :: i3_2_c, j3_2_c, k3_2_c
+ integer :: i3_3_s, j3_3_s, k3_3_s
+ integer :: i3_3_c, j3_3_c, k3_3_c
+ integer :: i3_4_s, j3_4_s, k3_4_s
+ integer :: i3_4_c, j3_4_c, k3_4_c
+ integer :: i3_5_s, j3_5_s, k3_5_s
+
+ !$acc parallel
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i0_1\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(i0_1\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i0_1 = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i0_2\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(j0_2\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(j0_2\\) firstprivate\\(i0_2\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i0_2 = 1, 100
+ do j0_2 = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+ do i1_s = 1, 100
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "gimple" } }
+ do i1_c = 1, 100
+ end do
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+ do i2_1_s = 1, 100
+ do j2_1_s = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+ do i2_1_c = 1, 100
+ do j2_1_c = 1, 100
+ end do
+ end do
+
+ !$acc parallel
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i2_2_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(i2_2_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i2_2_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+ do j2_2_s = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+ do i2_3_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+ do j2_3_s = 1, 100
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "gimple" } }
+ do i2_3_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "gimple" } }
+ do j2_3_c = 1, 100
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+ do i3_1_s = 1, 100
+ do j3_1_s = 1, 100
+ do k3_1_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+ do i3_1_c = 1, 100
+ do j3_1_c = 1, 100
+ do k3_1_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+ do i3_2_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+ do j3_2_s = 1, 100
+ do k3_2_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "gimple" } }
+ do i3_2_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "gimple" } }
+ do j3_2_c = 1, 100
+ do k3_2_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+ do i3_3_s = 1, 100
+ do j3_3_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+ do k3_3_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "gimple" } }
+ do i3_3_c = 1, 100
+ do j3_3_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "gimple" } }
+ do k3_3_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+ do i3_4_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+ do j3_4_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+ do k3_4_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc parallel loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "gimple" } }
+ do i3_4_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "gimple" } }
+ do j3_4_c = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "gimple" } }
+ do k3_4_c = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc parallel
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "firstprivate\\(i3_5_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(i3_5_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ do i3_5_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+ do j3_5_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+ do k3_5_s = 1, 100
+ end do
+ end do
+ end do
+ !$acc end parallel
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95
new file mode 100644
index 000000000000..6190669d63ae
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95
@@ -0,0 +1,142 @@
+! Predetermined 'private' clauses related to 'do' loops inside an OpenACC
+! accelerator routine.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+! The PR90114 XFAILs need to scan for the appropriate predetermined private
+! level.
+
+subroutine test
+ implicit none
+ integer :: i0_1
+ integer :: i0_2, j0_2
+ integer :: i1_s
+ integer :: i2_1_s, j2_1_s
+ integer :: i2_2_s, j2_2_s
+ integer :: i2_3_s, j2_3_s
+ integer :: i3_1_s, j3_1_s, k3_1_s
+ integer :: i3_2_s, j3_2_s, k3_2_s
+ integer :: i3_3_s, j3_3_s, k3_3_s
+ integer :: i3_4_s, j3_4_s, k3_4_s
+ integer :: i3_5_s, j3_5_s, k3_5_s
+ !$acc routine gang
+
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i0_1 = 1, 100
+ end do
+
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i0_2 = 1, 100
+ do j0_2 = 1, 100
+ end do
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+ do i1_s = 1, 100
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+ do i2_1_s = 1, 100
+ do j2_1_s = 1, 100
+ end do
+ end do
+
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i2_2_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+ do j2_2_s = 1, 100
+ end do
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+ do i2_3_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+ do j2_3_s = 1, 100
+ end do
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+ do i3_1_s = 1, 100
+ do j3_1_s = 1, 100
+ do k3_1_s = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+ do i3_2_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+ do j3_2_s = 1, 100
+ do k3_2_s = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+ do i3_3_s = 1, 100
+ do j3_3_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+ do k3_3_s = 1, 100
+ end do
+ end do
+ end do
+
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+ do i3_4_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+ do j3_4_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+ do k3_4_s = 1, 100
+ end do
+ end do
+ end do
+
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+ ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+ do i3_5_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+ do j3_5_s = 1, 100
+ !$acc loop independent
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+ do k3_5_s = 1, 100
+ end do
+ end do
+ end do
+end subroutine test
--
2.17.1
^ permalink raw reply [flat|nested] 12+ messages in thread
end of thread, other threads:[~2019-04-17 8:38 UTC | newest]
Thread overview: 12+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2014-08-11 23:55 [patch,gomp4] make fortran loop variables implicitly private in openacc Cesar Philippidis
2014-08-13 3:55 ` Cesar Philippidis
2014-08-13 20:41 ` Tobias Burnus
2014-08-14 11:15 ` Thomas Schwinge
2014-08-22 15:08 ` Cesar Philippidis
2014-09-08 6:26 ` Tobias Burnus
2014-09-09 1:25 ` Cesar Philippidis
2017-01-27 15:03 ` [gomp4] don't error on implicitly private induction variables in gfortran Cesar Philippidis
2017-01-27 15:07 ` Jakub Jelinek
2017-01-27 15:17 ` Cesar Philippidis
2018-08-07 22:04 ` [PATCH][OpenACC] Don't " Cesar Philippidis
2019-04-17 8:38 ` [PR90048] Fortran OpenACC 'private' clause rejected for predetermined private loop iteration variable (was: [patch,gomp4] make fortran loop variables implicitly private in openacc) 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).