From patchwork Thu Oct 23 07:41:57 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 402399 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 33E3F14007D for ; Thu, 23 Oct 2014 18:42:42 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=BJnT6H6Hu46/0tkr pBxmBFc0gJ781fTJxOOBzN7eyJjcr4jQPV0ga/Fi7cmusCkA5vJ7WypFXI5WlEEq ATgU7uV+m6NMXfWyGyozReoHY8WC2vvsP7nCvwPxkImEJWUw9M81hnJJ6DPG+ZHJ YOQbFpyNkjZ7ImDlOkNbqlPA/kw= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=V7rjb/mRLrhMGIGg8jPypZ JeVZI=; b=HPmJwt+0qY0pOMWTzNUnw8sKUA3LO+oyTWIrN1IJIUbqmIVaC92JRi qHXstXE0B8Je0mHf9WghWeC8oA3rnnWY+5KZwfNuWm03CWzMxp6lgp3waOzSsuMw nyodVrgP7U9jm+/1Gi+AinSUIQdCrYNB7iK2/RW2oyvWhvT61PvnU= Received: (qmail 26650 invoked by alias); 23 Oct 2014 07:42:35 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 26637 invoked by uid 89); 23 Oct 2014 07:42:34 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 23 Oct 2014 07:42:32 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1XhD2R-0002gh-Np from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 23 Oct 2014 00:42:28 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.181.6; Thu, 23 Oct 2014 08:42:26 +0100 From: Thomas Schwinge To: "gcc-patches@gcc.gnu.org" CC: Cesar Philippidis , Tom de Vries Subject: Re: [patch,gomp-4_0-branch] openacc parallel reduction part 1 In-Reply-To: <53D68A03.2050701@mentor.com> References: <53B9D780.4040904@mentor.com> <87wqaxh0x7.fsf@kepler.schwinge.homeip.net> <53D68A03.2050701@mentor.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (i586-pc-linux-gnu) Date: Thu, 23 Oct 2014 09:41:57 +0200 Message-ID: <87y4s7i662.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Mon, 28 Jul 2014 10:36:03 -0700, Cesar Philippidis wrote: > On 07/28/2014 10:02 AM, Thomas Schwinge wrote: > > On Sun, 6 Jul 2014 16:10:56 -0700, Cesar Philippidis 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 > > > > 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 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(-) Grüße, Thomas 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 + + * omp-low.c (finalize_reduction_data): Don't skip the reduction + variable's original value. + 2014-10-20 Cesar Philippidis * 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 + + * testsuite/libgomp.oacc-c/reduction-initial-1.c: New file. + 2014-10-20 Cesar Philippidis * (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; +}