From patchwork Mon Jul 28 17:02:12 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 374276 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 9E3F114013F for ; Tue, 29 Jul 2014 03:02:47 +1000 (EST) 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=slkVAYo7nAnK/2Ls WZNPcbNBBjDx53wpZTKoVfvp5nwWWxAvf+iF4yVmwKudsTxE7Ab6DTgQxXx3UhXa zYtbaazmG8IP28oDiwAlLqTP2QirF+vMuhmQrf95LMeARiJ1eidPV/4tURnVuIWP Dc5WRKBt0VTLDCxpskvAZ4oDKNY= 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=V0n1kMdMVnjycXVpNMvTuw 4FK7M=; b=Idp8mDNFoQlR0fei4Utw5t9RzxyRaL7DAU0KPKHGZwnKkMFfCezsDW 7zMhQUvCJM0VmwMrUqxMFAzHvh/nfaTG/hVDnLbtQGn03r+xpm0iKstqNXeMv5s8 dbScUZ4dNyfCm2PlQ4wFANUBodMDaNt7BVUsUCs/0AjHA8GVoSrX0= Received: (qmail 16713 invoked by alias); 28 Jul 2014 17:02:32 -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 16597 invoked by uid 89); 28 Jul 2014 17:02:31 -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 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; Mon, 28 Jul 2014 17:02:29 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1XBoJe-0007LR-4A from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Mon, 28 Jul 2014 10:02:26 -0700 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Mon, 28 Jul 2014 10:02:25 -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.2.247.3; Mon, 28 Jul 2014 18:02:23 +0100 From: Thomas Schwinge To: Cesar Philippidis CC: "gcc-patches@gcc.gnu.org" Subject: Re: [patch,gomp-4_0-branch] openacc parallel reduction part 1 In-Reply-To: <53B9D780.4040904@mentor.com> References: <53B9D780.4040904@mentor.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (i486-pc-linux-gnu) Date: Mon, 28 Jul 2014 19:02:12 +0200 Message-ID: <87wqaxh0x7.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi Cesar! 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? Grüße, Thomas --- 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. */