Hi Frederik! On 2019-10-29T13:20:53+0100, "Harwath, Frederik" <frede...@codesourcery.com> wrote: > On 24.10.19 16:31, Thomas Schwinge wrote: >> So just C/C++ testing, no Fortran at all. This is not ideal, but >> probably (hopefully) acceptable given that this is working on the middle >> end representation shared between all front ends. > > Thanks to Tobias, we now also have Fortran tests.
Indeed, thanks, Tobias. I have not reviewed these in great detail, but
they certainly do look plausible.
>>> --- a/gcc/testsuite/c-c++-common/goacc/reduction-6.c
>>> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
>>> @@ -16,17 +16,6 @@ int foo (int N)
>>> }
>>> }
>>>
>>> - #pragma acc parallel
>>> - {
>>> - #pragma acc loop reduction(+:b)
>>> - for (int i = 0; i < N; i++)
>>> - {
>>> - #pragma acc loop
>>> - for (int j = 0; j < N; j++)
>>> - b += 1;
>>> - }
>>> - }
>>> -
>>> #pragma acc parallel
>>> {
>>> #pragma acc loop reduction(+:c)
>>
>> That one stays in, but gets a 'dg-warning'.
>
> What warning would you expect to see here? I do not get any warnings.
What I meant was that you should re-instantiate the code removed here,
and then add the expected 'dg-warning'.
..., but upon having a look myself, I notice that there actually is no
"nested loop in reduction needs reduction clause" diagnostic printed
here, huh. Should there be? (OK to address separately, later on.)
Similar for the libgomp execution test cases: undo the 'reduction' clause
additions, and instead add the expected 'dg-warning's (here, they're
really necessary), for the reason I had given at the end of my email.
Sorry if that was unclear.
For the same reason, please also leave out Tobias' translated
'libgomp.oacc-fortran/par-loop-comb-reduction-1.f90' -- we shall later
consider that one, separately.
For your convenience, I'm attaching an incremental patch, to be merged
into yours.
> From 22f45d4c2c11febce171272f9289c487aed4f9d7 Mon Sep 17 00:00:00 2001
> From: Frederik Harwath <frede...@codesourcery.com>
> Date: Tue, 29 Oct 2019 12:39:23 +0100
> Subject: [PATCH] Warn about inconsistent OpenACC nested reduction clauses
> MIME-Version: 1.0
> Content-Type: text/plain; charset=UTF-8
> Content-Transfer-Encoding: 8bit
>
> OpenACC (cf. OpenACC 2.7, section 2.9.11. "reduction clause";
> this was first clarified by OpenACC 2.6) requires that, if a
> variable is used in reduction clauses on two nested loops, then
> there must be reduction clauses for that variable on all loops
> that are nested in between the two loops and all these reduction
> clauses must use the same operator.
> This commit introduces a check for that property which reports
> warnings if it is violated.
>
> In gcc/testsuite/c-c++-common/goacc/reduction-6.c, we remove the erroneous
> reductions on variable b; adding a reduction clause to make it compile
> cleanly
> would make it a duplicate of the test for variable c.
The latter paragraph then is not needed anymore.
> 2019-10-29 Gergö Barany <ge...@codesourcery.com>
> Tobias Burnus <tob...@codesourcery.com>
> Frederik Harwath <frede...@codesourcery.com>
> Thomas Schwinge <tho...@codesourcery.com>
>
> gcc/
> * omp-low.c (struct omp_context): New fields
> local_reduction_clauses, outer_reduction_clauses.
> (new_omp_context): Initialize these.
> (scan_sharing_clauses): Record reduction clauses on OpenACC constructs.
> (scan_omp_for): Check reduction clauses for incorrect nesting.
> gcc/testsuite/
> * c-c++-common/goacc/nested-reductions-warn.c: New test.
> * c-c++-common/goacc/nested-reductions.c: New test.
> * c-c++-common/goacc/reduction-6.c: Adjust.
> * gfortran.dg/goacc/nested-reductions-warn.f90: New test.
> * gfortran.dg/goacc/nested-reductions.f90: New test.
> libgomp/
> * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
> Add missing reduction clauses.
> * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
> Likewise.
> * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
> Likewise.
> * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
> Likewise.
> * testsuite/libgomp.oacc-fortran/par-loop-comb-reduction-1.f90:
> New test.
The ChangeLog updates still have to be adjusted per my incremental patch.
With that addressed, OK for trunk.
A few more comments to address separately, later on.
I noticed in the 'libgomp.log' that we currently print:
[...]/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c: In function
'main':
[...]/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:18:13: warning:
nested loop in reduction needs reduction clause for 'res'
[...]/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:18:13: warning:
nested loop in reduction needs reduction clause for 'res'
Duplicate diagnostic, due to the the two nested inner loops. (I'm just
noting that, not complaining.)
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -2410,6 +2425,89 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
> + ctx->outer_reduction_clauses
> + = chainon (unshare_expr (ctx->outer->local_reduction_clauses),
> + ctx->outer->outer_reduction_clauses);
I have not quickly re-understood -- if such a word exists ;-) -- the
logic here, but maybe we don't need anymore to 'unshare_expr', as we're
now no longer modifying the 'local_clause'...
> + tree outer_reduction_clauses = ctx->outer_reduction_clauses;
> + tree local_iter = local_reduction_clauses;
> + for (; local_iter; local_iter = TREE_CHAIN (local_iter))
> + {
> + tree local_clause = TREE_VALUE (local_iter);
> + tree local_var = OMP_CLAUSE_DECL (local_clause);
> + tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause);
> + bool have_outer_reduction = false;
> + tree ctx_iter = outer_reduction_clauses;
> + for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter))
> + {
> + tree outer_clause = TREE_VALUE (ctx_iter);
> + tree outer_var = OMP_CLAUSE_DECL (outer_clause);
> + tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause);
> + if (outer_var == local_var && outer_op != local_op)
> + {
> + warning_at (gimple_location (stmt), 0,
> + "conflicting reduction operations for %qE",
> + local_var);
> + inform (OMP_CLAUSE_LOCATION (outer_clause),
> + "location of the previous reduction for %qE",
> + outer_var);
> +
> + }
... here. (Once done to "suppress spurious errors".)
> + if (outer_var == local_var)
> + {
> + have_outer_reduction = true;
> + break;
> + }
> + }
> + if (have_outer_reduction)
> + {
> + /* There is a reduction on outer_var both on this loop and on
> + some enclosing loop. Walk up the context tree until such a
> + loop with a reduction on outer_var is found, and complain
> + about all intervening loops that do not have such a
> + reduction. */
> + struct omp_context *curr_loop = ctx->outer;
> + bool found = false;
> + while (curr_loop != NULL)
> + {
> + tree curr_iter = curr_loop->local_reduction_clauses;
> + for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter))
> + {
> + tree curr_clause = TREE_VALUE (curr_iter);
> + tree curr_var = OMP_CLAUSE_DECL (curr_clause);
> + if (curr_var == local_var)
> + {
> + found = true;
> + break;
> + }
> + }
> + if (!found)
> + warning_at (gimple_location (curr_loop->stmt), 0,
> + "nested loop in reduction needs "
> + "reduction clause for %qE",
> + local_var);
Might also be useful (if feasible?) to 'inform' what the outer/inner
'reduction' clauses are that trigger this.
> + else
> + break;
> + curr_loop = curr_loop->outer;
> + }
> + }
> + }
> + ctx->local_reduction_clauses = local_reduction_clauses;
> + ctx->outer_reduction_clauses
> + = chainon (unshare_expr (ctx->local_reduction_clauses),
> + ctx->outer_reduction_clauses);
> }
Same (?) comment as above about 'unshare_expr'.
Grüße
Thomas
From 7d96a9fa0d8dbddaa271f1bd22e6de3227009639 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Tue, 5 Nov 2019 15:12:13 +0100
Subject: [PATCH] into Warn about inconsistent OpenACC nested reduction clauses
---
gcc/omp-low.c | 1 -
.../c-c++-common/goacc/reduction-6.c | 11 ++++++
.../par-loop-comb-reduction-1.c | 2 +-
.../par-loop-comb-reduction-2.c | 2 +-
.../par-loop-comb-reduction-3.c | 2 +-
.../par-loop-comb-reduction-4.c | 2 +-
.../par-loop-comb-reduction-1.f90 | 35 -------------------
7 files changed, 15 insertions(+), 40 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-fortran/par-loop-comb-reduction-1.f90
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ebf95f1c649..f9d21655696 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2475,7 +2475,6 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
inform (OMP_CLAUSE_LOCATION (outer_clause),
"location of the previous reduction for %qE",
outer_var);
-
}
if (outer_var == local_var)
{
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-6.c b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
index 3c10b4dddaf..619f82b9d8b 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-6.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-6.c
@@ -16,6 +16,17 @@ int foo (int N)
}
}
+ #pragma acc parallel
+ {
+ #pragma acc loop reduction(+:b)
+ for (int i = 0; i < N; i++)
+ {
+ #pragma acc loop
+ for (int j = 0; j < N; j++)
+ b += 1;
+ }
+ }
+
#pragma acc parallel
{
#pragma acc loop reduction(+:c)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
index 91fe772045d..8c08717fdf2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c
@@ -15,7 +15,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) copy(res)
{
- #pragma acc loop gang reduction(+:res)
+ #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */
for (j = 0; j < 32; j++)
{
#pragma acc loop worker reduction(+:res)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
index 8a7bc724070..41042accbde 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c
@@ -14,7 +14,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(^:res)
{
- #pragma acc loop gang reduction(^:res)
+ #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(^:res)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
index eba5c6a544e..ace1e005e2e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
@@ -16,7 +16,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) copy(res)
{
- #pragma acc loop gang reduction(+:res)
+ #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(+:res)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
index 12b823f33ab..c3cc12fa953 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
@@ -16,7 +16,7 @@ main (int argc, char *argv[])
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) reduction(max:mres) copy(res, mres)
{
- #pragma acc loop gang reduction(+:res) reduction(max:mres)
+ #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'm\?res'" "TODO" } */
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(+:res) reduction(max:mres)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-loop-comb-reduction-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/par-loop-comb-reduction-1.f90
deleted file mode 100644
index 6fa684a6adf..00000000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/par-loop-comb-reduction-1.f90
+++ /dev/null
@@ -1,35 +0,0 @@
-! Test of reduction on both parallel and loop directives (worker and
-! vector-partitioned loops individually in gang-partitioned mode, integer
-! type).
-
-program main
- integer :: i, j, arr(0:32768-1), res, hres
- res = 0
- hres = 0
-
- arr = [(i, i = 1, 32768)]
-
- !$acc parallel num_gangs(32) num_workers(32) vector_length(32) &
- !$acc & reduction(+:res) copy(res)
- !$acc loop gang reduction(+:res)
- do j = 0, 31
- !$acc loop worker reduction(+:res)
- do i = 0, 1023
- res = res + arr(j * 1024 + i)
- end do
-
- !$acc loop vector reduction(+:res)
- do i = 1023, 0, -1
- res = res + arr(j * 1024 + i)
- end do
- end do
- !$acc end parallel
-
- do j = 0, 31
- do i = 0, 1023
- hres = hres + arr(j * 1024 + i) * 2
- end do
- end do
-
- if (res /= hres) stop 1
-end program main
--
2.17.1
signature.asc
Description: PGP signature
