Hi Cesar! On Sun, 6 Jul 2014 16:10:56 -0700, Cesar Philippidis <[email protected]> wrote: > This patch is the first step to enabling parallel reductions in openacc.
I think I have found one issue in this code -- but please verify that my
understanding of reductions is correct. Namely:
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> +/* Helper function to finalize local data for the reduction arrays. The
> + reduction array needs to be reduced to the original reduction variable.
> + FIXME: This function assumes that there are vector_length threads in
> + total. Also, it assumes that there are at least vector_length iterations
> + in the for loop. */
> +
> +static void
> +finalize_reduction_data (tree clauses, tree nthreads, gimple_seq *stmt_seqp,
> + omp_context *ctx)
> +{
> + gcc_assert (is_gimple_omp_oacc_specifically (ctx->stmt));
> +
> + tree c, var, array, loop_header, loop_body, loop_exit;
> + gimple stmt;
> +
> + /* Create for loop.
> +
> + let var = the original reduction variable
> + let array = reduction variable array
> +
> + var = array[0]
> + for (i = 1; i < nthreads; i++)
> + var op= array[i]
> + */
This should also consider the reduction variable's original value. Test
case (which does the expected thing if modified for OpenMP):
#include <stdlib.h>
int
main(void)
{
#define I 5
#define N 11
#define A 8
int a = A;
int s = I;
#pragma acc parallel vector_length(N)
{
int i;
#pragma acc loop reduction(+:s)
for (i = 0; i < N; ++i)
s += a;
}
if (s != I + N * A)
abort ();
return 0;
}
OK to check in the following?
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -9547,8 +9547,7 @@ finalize_reduction_data (tree clauses, tree nthreads,
gimple_seq *stmt_seqp,
let var = the original reduction variable
let array = reduction variable array
- var = array[0]
- for (i = 1; i < nthreads; i++)
+ for (i = 0; i < nthreads; i++)
var op= array[i]
*/
@@ -9556,42 +9555,9 @@ finalize_reduction_data (tree clauses, tree nthreads,
gimple_seq *stmt_seqp,
loop_body = create_artificial_label (UNKNOWN_LOCATION);
loop_exit = create_artificial_label (UNKNOWN_LOCATION);
- /* Initialize the reduction variables to be value of the first array
- element. */
- for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
- continue;
-
- tree_code reduction_code = OMP_CLAUSE_REDUCTION_CODE (c);
-
- /* reduction(-:var) sums up the partial results, so it acts
- identically to reduction(+:var). */
- if (reduction_code == MINUS_EXPR)
- reduction_code = PLUS_EXPR;
-
- /* Set up reduction variable, var. Becuase it's not gimple register,
- it needs to be treated as a reference. */
- var = OMP_CLAUSE_DECL (c);
- type = get_base_type (var);
- tree ptr = lookup_reduction (omp_get_id (OMP_CLAUSE_DECL (c)), ctx);
-
- /* Extract array[0] into mem. */
- tree mem = create_tmp_var (type, NULL);
- gimplify_assign (mem, build_simple_mem_ref (ptr), stmt_seqp);
-
- /* Find the original reduction variable. */
- tree x = build_outer_var_ref (var, ctx);
- if (is_reference (var))
- var = build_simple_mem_ref (var);
-
- x = lang_hooks.decls.omp_clause_assign_op (c, var, mem);
- gimplify_and_add (unshare_expr(x), stmt_seqp);
- }
-
- /* Create an index variable and set it to one. */
+ /* Create and initialize an index variable. */
tree ix = create_tmp_var (sizetype, NULL);
- gimplify_assign (ix, fold_build1 (NOP_EXPR, sizetype, integer_one_node),
+ gimplify_assign (ix, fold_build1 (NOP_EXPR, sizetype, integer_zero_node),
stmt_seqp);
/* Insert the loop header label here. */
Grüße,
Thomas
pgpurf6HU1CCO.pgp
Description: PGP signature
