+2018-12-20 Gergö Barany <gergo@codesourcery.com>
+
+ * omp-low.c (struct omp_context): New fields
+ local_reduction_clauses, outer_reduction_clauses.
+ (new_omp_context): Initialize these.
+ (scan_sharing_clauses): Record reduction clauses on OpenACC
+ constructs.
+ (scan_omp_for): Check reduction clauses for incorrect nesting.
+
2019-01-31 Kwok Cheung Yeung <kcy@codesourcery.com>
* omp-low.c (lower_omp_target): For use_device clauses, generate
/* Hash map of dynamic arrays in this context. */
hash_map<tree_operand_hash, tree> *dynamic_arrays;
+ /* A tree_list of the reduction clauses in this context. */
+ tree local_reduction_clauses;
+
+ /* A tree_list of the reduction clauses in outer contexts. */
+ tree outer_reduction_clauses;
+
/* The number of levels of OpenACC partitioning invoked in this context. */
unsigned oacc_partitioning_levels;
ctx->cb = outer_ctx->cb;
ctx->cb.block = NULL;
ctx->depth = outer_ctx->depth + 1;
+ ctx->local_reduction_clauses = NULL;
+ ctx->outer_reduction_clauses = ctx->outer_reduction_clauses;
}
else
{
ctx->cb.adjust_array_error_bounds = true;
ctx->cb.dont_remap_vla_if_no_change = true;
ctx->depth = 1;
+ ctx->local_reduction_clauses = NULL;
+ ctx->outer_reduction_clauses = NULL;
}
ctx->cb.decl_map = new hash_map<tree, tree>;
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
+ if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx))
+ ctx->local_reduction_clauses
+ = tree_cons (NULL, c, ctx->local_reduction_clauses);
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (decl) == MEM_REF)
{
gimple_omp_for_set_clauses (stmt, clauses);
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
+ a reduction on v somewhere in an outer context, then there is a
+ reduction on v on all intervening loops as well. */
+ tree local_reduction_clauses = NULL;
+ for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ local_reduction_clauses
+ = tree_cons (NULL, c, local_reduction_clauses);
+ }
+ if (ctx->outer_reduction_clauses == NULL && ctx->outer != NULL)
+ ctx->outer_reduction_clauses
+ = chainon (unshare_expr (ctx->outer->local_reduction_clauses),
+ ctx->outer->outer_reduction_clauses);
+ tree outer_reduction_clauses = ctx->outer_reduction_clauses;
+ tree local_iter = local_reduction_clauses;
+ for (; local_iter; local_iter = TREE_CHAIN (local_iter))
+ {
+ tree local_clause = TREE_VALUE (local_iter);
+ tree local_var = OMP_CLAUSE_DECL (local_clause);
+ tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause);
+ bool have_outer_reduction = false;
+ tree ctx_iter = outer_reduction_clauses;
+ for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter))
+ {
+ tree outer_clause = TREE_VALUE (ctx_iter);
+ tree outer_var = OMP_CLAUSE_DECL (outer_clause);
+ tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause);
+ if (outer_var == local_var && outer_op != local_op)
+ {
+ error_at (gimple_location (stmt),
+ "conflicting reduction operations for %qE",
+ local_var);
+ inform (OMP_CLAUSE_LOCATION (outer_clause),
+ "location of the previous reduction for %qE",
+ outer_var);
+ /* Change this operation to be equal to the outer one.
+ This is meant to suppress spurious errors; for example,
+ in nested +, -, + reductions, we would generate errors
+ for both the change from + to - and from - to +. */
+ OMP_CLAUSE_REDUCTION_CODE (local_clause) = outer_op;
+ /* Also change the location so that in nested +, -, -
+ reductions, the second error message also refers to the
+ outermost + reduction. */
+ OMP_CLAUSE_LOCATION (local_clause)
+ = OMP_CLAUSE_LOCATION (outer_clause);
+ }
+ if (outer_var == local_var)
+ {
+ have_outer_reduction = true;
+ break;
+ }
+ }
+ if (have_outer_reduction)
+ {
+ /* There is a reduction on outer_var both on this loop and on
+ some enclosing loop. Walk up the context tree until such a
+ loop with a reduction on outer_var is found, and complain
+ about all intervening loops that do not have such a
+ reduction. */
+ struct omp_context *curr_loop = ctx->outer;
+ bool found = false;
+ while (curr_loop != NULL)
+ {
+ tree curr_iter = curr_loop->local_reduction_clauses;
+ for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter))
+ {
+ tree curr_clause = TREE_VALUE (curr_iter);
+ tree curr_var = OMP_CLAUSE_DECL (curr_clause);
+ if (curr_var == local_var)
+ {
+ found = true;
+ break;
+ }
+ }
+ if (!found)
+ error_at (gimple_location (curr_loop->stmt),
+ "nested loop in reduction needs "
+ "reduction clause for %qE",
+ local_var);
+ else
+ break;
+ curr_loop = curr_loop->outer;
+ }
+ }
+ }
+ ctx->local_reduction_clauses = local_reduction_clauses;
+ ctx->outer_reduction_clauses
+ = chainon (unshare_expr (ctx->local_reduction_clauses),
+ ctx->outer_reduction_clauses);
}
scan_sharing_clauses (clauses, ctx);
+2018-12-20 Gergö Barany <gergo@codesourcery.com>
+
+ * c-c++-common/goacc/nested-reductions-fail.c: New test.
+ * c-c++-common/goacc/nested-reductions.c: New test.
+ * c-c++-common/goacc/reduction-6.c: Adjust.
+
2018-10-02 Cesar Philippidis <cesar@codesourcery.com>
* gfortran.dg/goacc/pr72741.f90: Update test for current diagnostics.
--- /dev/null
+/* 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;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "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-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ 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-error "conflicting reduction operations 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 reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "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-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "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-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "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-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;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "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 parallel loop construct. */
+
+void acc_parallel_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel loop
+ for (int h = 0; h < 10; ++h)
+ {
+ #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;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "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-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ 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-error "conflicting reduction operations 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 reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "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-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "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-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "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-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;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "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 now the outermost reduction clause is on
+ the parallel region, not the outermost loop. */
+void acc_parallel_reduction (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel reduction(+:sum)
+ {
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .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;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "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 // { dg-error "nested loop in reduction needs reduction clause for .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 // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { 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;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "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 parallel loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_parallel_loop_reduction (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel loop reduction(+:sum)
+ for (int h = 0; h < 10; ++h)
+ {
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .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;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "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 // { dg-error "nested loop in reduction needs reduction clause for .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 // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(+:sum)
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(max:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (i = 0; i < 10; i++)
+ #pragma acc loop reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (l = 0; l < 10; l++)
+ sum = 1;
+
+ #pragma acc loop reduction(-:diff) // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (i = 0; i < 10; i++)
+ {
+ #pragma acc loop reduction(-:diff) // { 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;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "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 inside a routine construct. */
+#pragma acc routine gang
+void acc_routine (void)
+{
+ int i, j, k, l, sum, diff;
+
+ {
+ #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;
+
+ #pragma acc loop reduction(+:sum)
+ for (i = 0; i < 10; i++)
+ #pragma acc loop collapse(2) // { dg-error "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-error "nested loop in reduction needs reduction clause for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ 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-error "conflicting reduction operations 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 reduction(-:sum) // { dg-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(-:sum) // { dg-error "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-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop // { dg-error "nested loop in reduction needs reduction clause for .sum." }
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "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-error "conflicting reduction operations for .sum." }
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ #pragma acc loop reduction(*:sum) // { dg-error "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-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;
+
+ #pragma acc loop reduction(+:sum) // { dg-error "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;
+ }
+ }
+}
+
+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;
+
+ #pragma acc loop reduction(+:sum)
+ 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++)
+ #pragma acc loop reduction(-:diff)
+ 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
+ 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;
+ }
+}
--- /dev/null
+/* Test cases of nested reduction loops that should compile cleanly. */
+
+void acc_parallel (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc parallel
+ {
+ #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)
+ 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 parallel loop construct. */
+
+void acc_parallel_loop (void)
+{
+ int i, j, k, l, sum, diff;
+
+ #pragma acc parallel 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" }
+ 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" }
+ 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" }
+ 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 parallel region, not the outermost loop. */
+
+void acc_parallel_reduction (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc parallel reduction(+:sum)
+ {
+ 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)
+ 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;
+ }
+
+ #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(-:diff)
+ 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)
+ for (j = 0; j < 10; j++)
+ #pragma acc loop reduction(+:sum)
+ for (k = 0; k < 10; k++)
+ sum = 1;
+
+ #pragma acc loop
+ 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 parallel loop construct, and
+ the outermost reduction clause is on that one, not the outermost loop. */
+void acc_parallel_loop_reduction (void)
+{
+ int i, j, k, sum, diff;
+
+ #pragma acc parallel 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" }
+ 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" }
+ 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" }
+ 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" }
+ 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" }
+ 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" }
+ 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" }
+ 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 inside a routine construct. */
+#pragma acc routine gang
+void acc_routine (void) // { dg-bogus "region is gang partitioned but does not contain gang partitioned code" "TODO" { xfail *-*-* } }
+{
+ int i, j, k, sum, diff;
+
+ {
+ #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-bogus "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-bogus "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-bogus "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;
+
+ /* 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 (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++)
+ 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 parallel
- {
- #pragma acc loop reduction(+:b)
- for (int i = 0; i < N; i++)
- {
- #pragma acc loop
- for (int j = 0; j < N; j++)
- b += 1;
- }
- }
-
#pragma acc parallel
{
#pragma acc loop reduction(+:c)
+2018-12-20 Gergö Barany <gergo@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
+ Add missing reduction clauses.
+ * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
+ Likewise.
+
2019-01-30 Kwok Cheung Yeung <kcy@codesourcery.com>
* testsuite/libgomp.oacc-fortran/optional-cache.f95
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) copy(res)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(+:res)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker reduction(+:res)
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(^:res)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(^:res)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(^:res)
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) copy(res)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(+:res)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(+:res)
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
reduction(+:res) reduction(max:mres) copy(res, mres)
{
- #pragma acc loop gang
+ #pragma acc loop gang reduction(+:res) reduction(max:mres)
for (j = 0; j < 32; j++)
{
#pragma acc loop worker vector reduction(+:res) reduction(max:mres)
projects / thirdparty / gcc.git / commitdiff
