Hi!
On Mon, 28 Jul 2014 10:36:03 -0700, Cesar Philippidis
<cesar_philippi...@mentor.com> wrote:
> On 07/28/2014 10:02 AM, Thomas Schwinge wrote:
> > On Sun, 6 Jul 2014 16:10:56 -0700, Cesar Philippidis
> > <cesar_philippi...@mentor.com> 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?
>
> Reductions can be specified with both the parallel and loop constructs.
> According to section 2.5.11 in the opacc spec, a reduction in a parallel
> construct should behave as you described:
>
> At the end of the region, the values for each gang are combined
> using the reduction operator, and the result combined with the
> value of the original variable and stored in the original
> variable.
>
> However,in section 2.7.11, a reduction in a loop construct behaves as
> follows:
>
> At the end of the loop, the values for each thread are combined
> using the specified reduction operator, and the result stored
> in the original variable at the end of the parallel or kernels
> region.
>
> The parallel reduction behavior does make more sense though. I'll ask
> the openacc gurus if there's a typo in section 2.7.11. It does refer to
> parallel reduction.
I proceeded by checking in the following patch to gomp-4_0-branch,
r216574:
commit 75e2a58b8ef7d20be2239ff029493986542ee7e3
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Oct 23 07:26:40 2014 +0000
OpenACC reductions: Don't skip the reduction variable's original value.
gcc/
* omp-low.c (finalize_reduction_data): Don't skip the reduction
variable's original value.
libgomp/
* testsuite/libgomp.oacc-c/reduction-initial-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@216574
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 5 +++
gcc/omp-low.c | 40 ++--------------------
libgomp/ChangeLog.gomp | 4 +++
.../testsuite/libgomp.oacc-c/reduction-initial-1.c | 32 +++++++++++++++++
4 files changed, 44 insertions(+), 37 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 6d107d2..28e7252 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-10-23 Thomas Schwinge <tho...@codesourcery.com>
+
+ * omp-low.c (finalize_reduction_data): Don't skip the reduction
+ variable's original value.
+
2014-10-20 Cesar Philippidis <ce...@codesourcery.com>
* gimplify.c (gimplify_scan_omp_clauses): Remove switch stmt which
diff --git gcc/omp-low.c gcc/omp-low.c
index b8022c2..b21235f 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -9869,8 +9869,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]
*/
@@ -9878,42 +9877,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. */
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 065dfb6..5363068 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2014-10-23 Thomas Schwinge <tho...@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c/reduction-initial-1.c: New file.
+
2014-10-20 Cesar Philippidis <ce...@codesourcery.com>
* (GOACC_update): Declare.
diff --git libgomp/testsuite/libgomp.oacc-c/reduction-initial-1.c
libgomp/testsuite/libgomp.oacc-c/reduction-initial-1.c
new file mode 100644
index 0000000..e763cf2
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/reduction-initial-1.c
@@ -0,0 +1,32 @@
+/* { dg-do run } */
+/* TODO:
+ { dg-xfail-run-if "" { *-*-* } { "-DACC_DEVICE_TYPE_host=1" } { "" } } */
+
+int
+main(void)
+{
+#define I 5
+/* TODO */
+#ifdef ACC_DEVICE_TYPE_host_nonshm
+# define N 1
+#else
+# define N 11
+#endif
+#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)
+ __builtin_abort();
+
+ return 0;
+}
Grüße,
Thomas
pgpXG_Kis7ElU.pgp
Description: PGP signature
