Hi!
On 2020-11-03T09:39:54+0100, I wrote:
> On 2018-12-20T15:28:33+0100, I wrote:
>> On behalf of Gergő (who doesn't have write access yet) I've pushed the
>> attached to openacc-gcc-8-branch.
>
> (Which then eventually got into master branch via Frederik.)
>> --- /dev/null
>> +++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-fail.c
>> @@ -0,0 +1,492 @@
>> +/* Test erroneous cases of nested reduction loops. */
>> +
>> +void acc_parallel (void)
>> +{
>> + int i, j, k, l, sum, diff;
>> +
>> + #pragma acc parallel
>> + {
>> + #pragma acc loop reduction(+:sum)
>> + for (i = 0; i < 10; i++)
>> + #pragma acc loop // { dg-error "nested loop in reduction needs
>> reduction clause for .sum." }
>> + for (j = 0; j < 10; j++)
>> + #pragma acc loop reduction(+:sum)
>> + for (k = 0; k < 10; k++)
>> + sum = 1;
>
>> +void acc_kernels (void)
>> +{
>> + int i, j, k, sum, diff;
>> +
>> + /* FIXME: No diagnostics are produced for these loops because reductions
>> + in kernels regions are not supported yet. */
>> + #pragma acc kernels
>> + {
>> + #pragma acc loop reduction(+:sum)
>> + for (i = 0; i < 10; i++)
>> + for (j = 0; j < 10; j++)
>> + for (k = 0; k < 10; k++)
>> + sum = 1;
Getting these diagnostics consistent is easy enough, however. I've
pushed "[OpenACC] Enable inconsistent nested 'reduction' clauses checking
for OpenACC 'kernels'" to master branch in commit
64dc14b1a764bd3059170431c9b43c6192dbd48f, and backported to
releases/gcc-10 branch in commit
217fb4d4e59e7d6e03a3704f80f401e2a641dbe5, see attached.
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
>From 64dc14b1a764bd3059170431c9b43c6192dbd48f Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <[email protected]>
Date: Thu, 22 Oct 2020 11:04:22 +0200
Subject: [PATCH] [OpenACC] Enable inconsistent nested 'reduction' clauses
checking for OpenACC 'kernels'
gcc/
* omp-low.c (scan_omp_for) <OpenACC>: Move earlier inconsistent
nested 'reduction' clauses checking.
gcc/testsuite/
* c-c++-common/goacc/nested-reductions-1-kernels.c: Extend.
* c-c++-common/goacc/nested-reductions-2-kernels.c: Likewise.
* gfortran.dg/goacc/nested-reductions-1-kernels.f90: Likewise.
* gfortran.dg/goacc/nested-reductions-2-kernels.f90: Likewise.
---
gcc/omp-low.c | 36 +-
.../goacc/nested-reductions-1-kernels.c | 199 +++++++++-
.../goacc/nested-reductions-2-kernels.c | 271 +++++++++++++-
.../goacc/nested-reductions-1-kernels.f90 | 251 ++++++++++++-
.../goacc/nested-reductions-2-kernels.f90 | 346 +++++++++++++++++-
5 files changed, 1063 insertions(+), 40 deletions(-)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index de5142f979b0..2f1a544bd46e 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2454,23 +2454,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
}
if (tgt && is_oacc_kernels (tgt))
- {
- /* Strip out reductions, as they are not handled yet. */
- tree *prev_ptr = &clauses;
-
- while (tree probe = *prev_ptr)
- {
- tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
-
- if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
- *prev_ptr = *next_ptr;
- else
- prev_ptr = next_ptr;
- }
-
- gimple_omp_for_set_clauses (stmt, clauses);
- check_oacc_kernel_gwv (stmt, ctx);
- }
+ check_oacc_kernel_gwv (stmt, ctx);
/* Collect all variables named in reductions on this loop. Ensure
that, if this loop has a reduction on some variable v, and there is
@@ -2553,6 +2537,24 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
ctx->outer_reduction_clauses
= chainon (unshare_expr (ctx->local_reduction_clauses),
ctx->outer_reduction_clauses);
+
+ if (tgt && is_oacc_kernels (tgt))
+ {
+ /* Strip out reductions, as they are not handled yet. */
+ tree *prev_ptr = &clauses;
+
+ while (tree probe = *prev_ptr)
+ {
+ tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
+
+ if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
+ *prev_ptr = *next_ptr;
+ else
+ prev_ptr = next_ptr;
+ }
+
+ gimple_omp_for_set_clauses (stmt, clauses);
+ }
}
scan_sharing_clauses (clauses, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
index 68cb8f82ee57..9323e2c8d7e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
@@ -6,8 +6,6 @@ void acc_kernels (void)
{
int i, j, k, sum, diff;
- /* FIXME: These tests are not meaningful yet because reductions in
- kernels regions are not supported yet. */
#pragma acc kernels
{
#pragma acc loop reduction(+:sum)
@@ -16,6 +14,12 @@ void acc_kernels (void)
for (k = 0; k < 10; k++)
sum = 1;
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
@@ -25,17 +29,208 @@ void acc_kernels (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined kernels loop construct. */
+
+void acc_kernels_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc kernels loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+ the kernels region, not the outermost loop. */
+
+void acc_kernels_reduction (void)
+{
+ /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+ supported on the 'kernels' construct. */
+}
+
+/* The same tests as above, but using a combined kernels loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_kernels_loop_reduction (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc kernels loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
}
}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
index 80d7c53daca7..dec7dbda45dc 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
@@ -4,36 +4,107 @@
void acc_kernels (void)
{
- int i, j, k, sum, diff;
+ int i, j, k, l, sum, diff;
- /* FIXME: No diagnostics are produced for these loops because reductions
- in kernels regions are not supported yet. */
#pragma acc kernels
{
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop
+ #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop reduction(-:diff)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
for (j = 0; j < 10; j++)
- #pragma acc loop
+ #pragma acc loop reduction(-:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined kernels loop construct. */
+
+void acc_kernels_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc kernels loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
@@ -41,10 +112,194 @@ void acc_kernels (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop reduction(-:sum)
+ #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+ the kernels region, not the outermost loop. */
+void acc_kernels_reduction (void)
+{
+ /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+ supported on the 'kernels' construct. */
+}
+
+/* The same tests as above, but using a combined kernels loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_kernels_loop_reduction (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc kernels loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." })
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
}
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
index 0999f8767c6e..60cb63006c3c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
@@ -6,8 +6,6 @@ subroutine acc_kernels ()
implicit none (type, external)
integer :: i, j, k, sum, diff
- ! FIXME: These tests are not meaningful yet because reductions in
- ! kernels regions are not supported yet.
!$acc kernels
!$acc loop reduction(+:sum)
do i = 1, 10
@@ -18,6 +16,15 @@ subroutine acc_kernels ()
end do
end do
+ !$acc loop collapse(2) reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
@@ -28,6 +35,15 @@ subroutine acc_kernels ()
end do
end do
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop collapse(2) reduction(+:sum)
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
!$acc loop reduction(+:sum)
do i = 1, 10
@@ -39,7 +55,6 @@ subroutine acc_kernels ()
end do
end do
-
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
@@ -51,5 +66,235 @@ subroutine acc_kernels ()
end do
end do
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(+:sum)
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff)
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
!$acc end kernels
end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+ implicit none (type, external)
+ integer :: h, i, j, k, l, sum, diff
+
+ !$acc kernels loop
+ do h = 1, 10
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop collapse(2) reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum)
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop collapse(2) reduction(+:sum)
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop. */
+
+subroutine acc_kernels_reduction ()
+ implicit none (type, external)
+
+ ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+ ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! The same tests as above, but using a combined kernels loop construct, and
+! the outermost reduction clause is on that one, not the outermost loop. */
+subroutine acc_kernels_loop_reduction ()
+ implicit none (type, external)
+ integer :: h, i, j, k, sum, diff
+
+ !$acc kernels loop reduction(+:sum)
+ do h = 1, 10
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ do i = 1, 10
+ !$acc loop
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ do i = 1, 10
+ do j = 1, 10
+ !$acc loop
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop_reduction
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
index edfd8625faf3..6ee41843d31a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
@@ -3,14 +3,15 @@
! See also 'c-c++-common/goacc/nested-reductions-2-kernels.c'.
subroutine acc_kernels ()
- integer :: i, j, k, sum, diff
+ implicit none (type, external)
+ integer :: i, j, k, l, sum, diff
- ! FIXME: No diagnostics are produced for these loops because reductions
- ! in kernels regions are not supported yet.
!$acc kernels
!$acc loop reduction(+:sum)
do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
do j = 1, 10
+ !$acc loop reduction(+:sum)
do k = 1, 10
sum = 1
end do
@@ -19,19 +20,37 @@ subroutine acc_kernels ()
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop
+ !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
do j = 1, 10
do k = 1, 10
- sum = 1
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
end do
end do
end do
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop reduction(-:diff)
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do j = 1, 10
- !$acc loop
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do k = 1, 10
sum = 1
end do
@@ -40,9 +59,9 @@ subroutine acc_kernels ()
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do j = 1, 10
- !$acc loop reduction(+:sum)
+ !$acc loop reduction(-:sum)
do k = 1, 10
sum = 1
end do
@@ -51,13 +70,320 @@ subroutine acc_kernels ()
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop reduction(-:sum)
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
sum = 1
end do
end do
+
+ !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
end do
!$acc end kernels
end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+ implicit none (type, external)
+ integer :: h, i, j, k, l, sum, diff
+
+ !$acc kernels loop
+ do h = 1, 10
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop.
+
+subroutine acc_kernels_reduction ()
+ implicit none (type, external)
+
+ ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+ ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! The same tests as above, but using a combined kernels loop construct, and
+! the outermost reduction clause is on that one, not the outermost loop. */
+subroutine acc_kernels_loop_reduction ()
+ implicit none (type, external)
+ integer :: h, i, j, k, l, sum, diff
+
+ !$acc kernels loop reduction(+:sum)
+ do h = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop_reduction
--
2.17.1
>From 217fb4d4e59e7d6e03a3704f80f401e2a641dbe5 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <[email protected]>
Date: Thu, 22 Oct 2020 11:04:22 +0200
Subject: [PATCH] [OpenACC] Enable inconsistent nested 'reduction' clauses
checking for OpenACC 'kernels'
gcc/
* omp-low.c (scan_omp_for) <OpenACC>: Move earlier inconsistent
nested 'reduction' clauses checking.
gcc/testsuite/
* c-c++-common/goacc/nested-reductions-1-kernels.c: Extend.
* c-c++-common/goacc/nested-reductions-2-kernels.c: Likewise.
* gfortran.dg/goacc/nested-reductions-1-kernels.f90: Likewise.
* gfortran.dg/goacc/nested-reductions-2-kernels.f90: Likewise.
(cherry picked from commit 64dc14b1a764bd3059170431c9b43c6192dbd48f)
---
gcc/omp-low.c | 36 +-
.../goacc/nested-reductions-1-kernels.c | 199 +++++++++-
.../goacc/nested-reductions-2-kernels.c | 271 +++++++++++++-
.../goacc/nested-reductions-1-kernels.f90 | 251 ++++++++++++-
.../goacc/nested-reductions-2-kernels.f90 | 346 +++++++++++++++++-
5 files changed, 1063 insertions(+), 40 deletions(-)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 16afc3720af2..91e1bcc6b69a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2447,23 +2447,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
}
if (tgt && is_oacc_kernels (tgt))
- {
- /* Strip out reductions, as they are not handled yet. */
- tree *prev_ptr = &clauses;
-
- while (tree probe = *prev_ptr)
- {
- tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
-
- if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
- *prev_ptr = *next_ptr;
- else
- prev_ptr = next_ptr;
- }
-
- gimple_omp_for_set_clauses (stmt, clauses);
- check_oacc_kernel_gwv (stmt, ctx);
- }
+ check_oacc_kernel_gwv (stmt, ctx);
/* Collect all variables named in reductions on this loop. Ensure
that, if this loop has a reduction on some variable v, and there is
@@ -2546,6 +2530,24 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
ctx->outer_reduction_clauses
= chainon (unshare_expr (ctx->local_reduction_clauses),
ctx->outer_reduction_clauses);
+
+ if (tgt && is_oacc_kernels (tgt))
+ {
+ /* Strip out reductions, as they are not handled yet. */
+ tree *prev_ptr = &clauses;
+
+ while (tree probe = *prev_ptr)
+ {
+ tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
+
+ if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
+ *prev_ptr = *next_ptr;
+ else
+ prev_ptr = next_ptr;
+ }
+
+ gimple_omp_for_set_clauses (stmt, clauses);
+ }
}
scan_sharing_clauses (clauses, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
index 68cb8f82ee57..9323e2c8d7e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-1-kernels.c
@@ -6,8 +6,6 @@ void acc_kernels (void)
{
int i, j, k, sum, diff;
- /* FIXME: These tests are not meaningful yet because reductions in
- kernels regions are not supported yet. */
#pragma acc kernels
{
#pragma acc loop reduction(+:sum)
@@ -16,6 +14,12 @@ void acc_kernels (void)
for (k = 0; k < 10; k++)
sum = 1;
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
@@ -25,17 +29,208 @@ void acc_kernels (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum)
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined kernels loop construct. */
+
+void acc_kernels_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc kernels loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
#pragma acc loop reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) reduction(+:sum)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+ the kernels region, not the outermost loop. */
+
+void acc_kernels_reduction (void)
+{
+ /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+ supported on the 'kernels' construct. */
+}
+
+/* The same tests as above, but using a combined kernels loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_kernels_loop_reduction (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc kernels loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ #pragma acc loop
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ for (i = 0; i < 10; i++)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(+:sum) // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
}
}
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
index 80d7c53daca7..dec7dbda45dc 100644
--- a/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/nested-reductions-2-kernels.c
@@ -4,36 +4,107 @@
void acc_kernels (void)
{
- int i, j, k, sum, diff;
+ int i, j, k, l, sum, diff;
- /* FIXME: No diagnostics are produced for these loops because reductions
- in kernels regions are not supported yet. */
#pragma acc kernels
{
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop
+ #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop reduction(-:diff)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
for (j = 0; j < 10; j++)
- #pragma acc loop
+ #pragma acc loop reduction(-:sum)
for (k = 0; k < 10; k++)
sum = 1;
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but using a combined kernels loop construct. */
+
+void acc_kernels_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc kernels loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
@@ -41,10 +112,194 @@ void acc_kernels (void)
#pragma acc loop reduction(+:sum)
for (i = 0; i < 10; i++)
- #pragma acc loop reduction(-:sum)
+ #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) reduction(-:diff)
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
+ }
+}
+
+/* The same tests as above, but now the outermost reduction clause is on
+ the kernels region, not the outermost loop. */
+void acc_kernels_reduction (void)
+{
+ /* In contrast to the 'parallel' construct, the 'reduction' clause is not
+ supported on the 'kernels' construct. */
+}
+
+/* The same tests as above, but using a combined kernels loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_kernels_loop_reduction (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc kernels loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
for (j = 0; j < 10; j++)
#pragma acc loop reduction(+:sum)
for (k = 0; k < 10; k++)
sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum) // { dg-warning "conflicting reduction operations for .sum." })
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-warning "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(+:sum) // { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ // { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:diff)
+ for (k = 0; k < 10; k++)
+ diff = 1;
+ }
}
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
index 0999f8767c6e..60cb63006c3c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-1-kernels.f90
@@ -6,8 +6,6 @@ subroutine acc_kernels ()
implicit none (type, external)
integer :: i, j, k, sum, diff
- ! FIXME: These tests are not meaningful yet because reductions in
- ! kernels regions are not supported yet.
!$acc kernels
!$acc loop reduction(+:sum)
do i = 1, 10
@@ -18,6 +16,15 @@ subroutine acc_kernels ()
end do
end do
+ !$acc loop collapse(2) reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
@@ -28,6 +35,15 @@ subroutine acc_kernels ()
end do
end do
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop collapse(2) reduction(+:sum)
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
!$acc loop reduction(+:sum)
do i = 1, 10
@@ -39,7 +55,6 @@ subroutine acc_kernels ()
end do
end do
-
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
@@ -51,5 +66,235 @@ subroutine acc_kernels ()
end do
end do
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(+:sum)
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff)
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
!$acc end kernels
end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+ implicit none (type, external)
+ integer :: h, i, j, k, l, sum, diff
+
+ !$acc kernels loop
+ do h = 1, 10
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop collapse(2) reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum)
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop collapse(2) reduction(+:sum)
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop. */
+
+subroutine acc_kernels_reduction ()
+ implicit none (type, external)
+
+ ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+ ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! The same tests as above, but using a combined kernels loop construct, and
+! the outermost reduction clause is on that one, not the outermost loop. */
+subroutine acc_kernels_loop_reduction ()
+ implicit none (type, external)
+ integer :: h, i, j, k, sum, diff
+
+ !$acc kernels loop reduction(+:sum)
+ do h = 1, 10
+ do i = 1, 10
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ do i = 1, 10
+ !$acc loop
+ do j = 1, 10
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ do i = 1, 10
+ do j = 1, 10
+ !$acc loop
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop_reduction
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
index edfd8625faf3..6ee41843d31a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-reductions-2-kernels.f90
@@ -3,14 +3,15 @@
! See also 'c-c++-common/goacc/nested-reductions-2-kernels.c'.
subroutine acc_kernels ()
- integer :: i, j, k, sum, diff
+ implicit none (type, external)
+ integer :: i, j, k, l, sum, diff
- ! FIXME: No diagnostics are produced for these loops because reductions
- ! in kernels regions are not supported yet.
!$acc kernels
!$acc loop reduction(+:sum)
do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
do j = 1, 10
+ !$acc loop reduction(+:sum)
do k = 1, 10
sum = 1
end do
@@ -19,19 +20,37 @@ subroutine acc_kernels ()
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop
+ !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
do j = 1, 10
do k = 1, 10
- sum = 1
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
end do
end do
end do
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop reduction(-:diff)
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do j = 1, 10
- !$acc loop
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do k = 1, 10
sum = 1
end do
@@ -40,9 +59,9 @@ subroutine acc_kernels ()
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
do j = 1, 10
- !$acc loop reduction(+:sum)
+ !$acc loop reduction(-:sum)
do k = 1, 10
sum = 1
end do
@@ -51,13 +70,320 @@ subroutine acc_kernels ()
!$acc loop reduction(+:sum)
do i = 1, 10
- !$acc loop reduction(-:sum)
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
do j = 1, 10
!$acc loop reduction(+:sum)
do k = 1, 10
sum = 1
end do
end do
+
+ !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
end do
!$acc end kernels
end subroutine acc_kernels
+
+! The same tests as above, but using a combined kernels loop construct.
+
+subroutine acc_kernels_loop ()
+ implicit none (type, external)
+ integer :: h, i, j, k, l, sum, diff
+
+ !$acc kernels loop
+ do h = 1, 10
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum)
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) reduction(-:diff)
+ do i = 1, 10
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop
+
+! The same tests as above, but now the outermost reduction clause is on
+! the kernels region, not the outermost loop.
+
+subroutine acc_kernels_reduction ()
+ implicit none (type, external)
+
+ ! In contrast to the 'parallel' construct, the 'reduction' clause is not
+ ! supported on the 'kernels' construct.
+end subroutine acc_kernels_reduction
+
+! The same tests as above, but using a combined kernels loop construct, and
+! the outermost reduction clause is on that one, not the outermost loop. */
+subroutine acc_kernels_loop_reduction ()
+ implicit none (type, external)
+ integer :: h, i, j, k, l, sum, diff
+
+ !$acc kernels loop reduction(+:sum)
+ do h = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop collapse(2) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(+:sum)
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(max:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do k = 1, 10
+ !$acc loop reduction(*:sum) ! { dg-warning "conflicting reduction operations for .sum." }
+ do l = 1, 10
+ sum = 1
+ end do
+ end do
+ end do
+ end do
+
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ do i = 1, 10
+ !$acc loop reduction(-:diff) ! { dg-warning "nested loop in reduction needs reduction clause for .sum." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(+:sum)
+ do k = 1, 10
+ sum = 1
+ end do
+ end do
+
+ !$acc loop reduction(+:sum) ! { dg-warning "nested loop in reduction needs reduction clause for .diff." }
+ ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 }
+ do j = 1, 10
+ !$acc loop reduction(-:diff)
+ do k = 1, 10
+ diff = 1
+ end do
+ end do
+ end do
+ end do
+end subroutine acc_kernels_loop_reduction
--
2.17.1
