From 64dc14b1a764bd3059170431c9b43c6192dbd48f Mon Sep 17 00:00:00 2001 From: Thomas Schwinge 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) : 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 de5142f979b..2f1a544bd46 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 68cb8f82ee5..9323e2c8d7e 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 80d7c53daca..dec7dbda45d 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 0999f8767c6..60cb63006c3 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 edfd8625faf..6ee41843d31 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.30.2