}
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
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);
{
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)
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)
#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;
+ }
}
}
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++)
#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;
+ }
}
}
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
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)
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
end do
end do
-
!$acc loop reduction(+:sum)
do i = 1, 10
!$acc loop reduction(+:sum)
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
! 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
!$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
!$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
!$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