From patchwork Wed Nov 12 17:54:40 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 410068 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 EC3AD1400E7 for ; Thu, 13 Nov 2014 04:55:04 +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:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=Lp7dK7Y8mMFSkgVzB4G5X/dH9dLAf nUwezvKMFTDOzsyK38/p2kW/RZf0lWfr7MznsaWcT5uphXkVGynaAv7t5uG2RA4y laFVTIJ9CL2XkWi4liIz1si1eSi5j3G6CTRr3YW/e/VkOHVQr7IHxyv23EdOXUiw 9akytSQFdHBdms= 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:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=RxsEIoZdkZTvG2/o82u4ps7slzw=; b=GmJ JFmaEqAckkCZpIvugRTs5Ui/hxF6thXirQ0mk7WopwNF5vFrHjO4WNw1bgWehwQv lMfZuvNRvBTIqIbLieS0QSzhutSNLTms9dbA82ANszhsiN3HFU2CkGlLaVBeCz5h kcG3D3KOkC5aL9CVVkU0xeFAj5B4KJdmPqlmq4tY= Received: (qmail 8358 invoked by alias); 12 Nov 2014 17:54:57 -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 8346 invoked by uid 89); 12 Nov 2014 17:54:56 -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; Wed, 12 Nov 2014 17:54:54 +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 1Xoc80-0006Ag-Ho from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Wed, 12 Nov 2014 09:54:49 -0800 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; Wed, 12 Nov 2014 17:54:47 +0000 From: Thomas Schwinge To: Subject: Re: [PATCH] [gomp4] Initial support of OpenACC loop directive in C front-end. In-Reply-To: <53E0F6CA.1040701@codesourcery.com> References: <53283E04.6010501@samsung.com> <87ha6vipjf.fsf@schwinge.name> <87zjklkk2f.fsf@kepler.schwinge.homeip.net> <87mwbsh6t7.fsf@kepler.schwinge.homeip.net> <53E0F6CA.1040701@codesourcery.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (i586-pc-linux-gnu) Date: Wed, 12 Nov 2014 18:54:40 +0100 Message-ID: <87lhngcnj3.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Tue, 5 Aug 2014 08:22:50 -0700, Cesar Philippidis wrote: > On 07/29/2014 02:07 AM, Thomas Schwinge wrote: > > > On Thu, 20 Mar 2014 15:42:48 +0100, I wrote: > >> On Tue, 18 Mar 2014 14:50:44 +0100, I wrote: > >>> On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov wrote: > >>>> This patch introduces support of OpenACC loop directive (and combined > >>>> directives) in C front-end up to GENERIC. Currently no clause is allowed. > >>> > >>> Thanks! I had worked on a simpler patch, not yet dealing with combined > >>> clauses. Also, I have some work for the GIMPLE level, namely building on > >>> GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP. I'll post this > >>> soon. > >> > >> Here are the patches, committed in r208702..4 to gomp-4_0-branch. > > > > Cesar, I hope I'm not confusing things here, but I remember that you once > > pointed out that in Fortran OpenACC, we have to explicitly specify the > > loop iteration variable in a private clause, whereas the OpenACC > > specification says this needs to happen automatically (predetermined data > > attribute). > > In light of this, please also review whether the following > > gimplify_omp_for changes of mine (when I originally added the OACC_LOOP > > support) are correct. Evidently, this code still has TODO markers in it; > > this was well before you added preliminary support for the private > > clause. That is, should we use GOVD_PRIVATE also for OACC_LOOP? > > OMP has both a private and a shared clause. And those neither of those > clauses can share variables, i.e. an index variable cannot be both > private and shared. By default, index variables are private in OMP, so > the gimplification pass makes those variables explicitly private so that > later passes can check for errors. > > Why does this nevertheless currently work for C for loops without any > > private clause for the loop iteration variable? Maybe because in C, the > > loop iteration variable ends up in a register and that is not shared > > between threads? I do see a private clause being added during > > gimplification for the OpenMP C test case, but not for OpenACC -- which I > > assume is precisely due to the code I'm quoting below. > > OMP lowering pass makes the index variables local to the parallel > function/kernel. So each thread effectively has its private copy of the > index variable. I have now committed the following to gomp-4_0-branch in r217433: commit 8562c72ad30e73dbd77b9d4170f71a395357aa33 Author: tschwinge Date: Wed Nov 12 17:49:59 2014 +0000 OpenACC loop construct: predetermined data attributes for loop variables. gcc/ * gimplify.c (gimplify_omp_for): Eliminate govd_private; always use GOVD_PRIVATE. gcc/testsuite/ * c-c++-common/goacc/loop-private-1.c: New file. libgomp/testsuite/ * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file. * testsuite/libgomp.oacc-c/collapse-4.c: Turn into an execution test, check result. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Update. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217433 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 3 +++ gcc/gimplify.c | 13 ++++--------- gcc/testsuite/ChangeLog.gomp | 4 ++++ gcc/testsuite/c-c++-common/goacc/loop-private-1.c | 14 ++++++++++++++ libgomp/ChangeLog.gomp | 8 ++++++++ .../collapse-4.c | 12 +++++++----- libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c | 20 ++++++++++---------- libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c | 20 ++++++++++---------- libgomp/testsuite/libgomp.oacc-c/collapse-4.c | 6 ++++-- 9 files changed, 64 insertions(+), 36 deletions(-) Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 2008542..0269746 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-11-12 Thomas Schwinge + * gimplify.c (gimplify_omp_for): Eliminate govd_private; always + use GOVD_PRIVATE. + * omp-low.c (scan_oacc_offload, expand_oacc_offload) (lower_oacc_offload): Merge into scan_omp_target, expand_omp_target, lower_omp_target, respectively. Update all diff --git gcc/gimplify.c gcc/gimplify.c index 5aeb726..233ac56 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -6843,7 +6843,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) gimple_seq for_body, for_pre_body; int i; bool simd; - enum gimplify_omp_var_data govd_private; enum omp_region_type ort; bitmap has_decl_expr = NULL; @@ -6855,18 +6854,15 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) case CILK_FOR: case OMP_DISTRIBUTE: simd = false; - govd_private = GOVD_PRIVATE; ort = ORT_WORKSHARE; break; case OMP_SIMD: case CILK_SIMD: simd = true; - govd_private = GOVD_PRIVATE; ort = ORT_SIMD; break; case OACC_LOOP: simd = false; - govd_private = /* TODO */ GOVD_LOCAL; ort = /* TODO */ ORT_WORKSHARE; break; default: @@ -6928,7 +6924,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl)) || POINTER_TYPE_P (TREE_TYPE (decl))); - /* Make sure the iteration variable is some kind of private. */ + /* Make sure the iteration variable is private. */ tree c = NULL_TREE; tree c2 = NULL_TREE; if (orig_for_stmt != for_stmt) @@ -6957,7 +6953,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) } else { - gcc_assert (govd_private == GOVD_PRIVATE); bool lastprivate = (!has_decl_expr || !bitmap_bit_p (has_decl_expr, DECL_UID (decl))); @@ -7001,7 +6996,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) else if (omp_is_private (gimplify_omp_ctxp, decl, 0)) omp_notice_variable (gimplify_omp_ctxp, decl, true); else - omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN); + omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN); /* If DECL is not a gimple register, create a temporary variable to act as an iteration counter. This is valid, since DECL cannot be @@ -7036,7 +7031,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) } else omp_add_variable (gimplify_omp_ctxp, var, - govd_private | GOVD_SEEN); + GOVD_PRIVATE | GOVD_SEEN); } else var = decl; @@ -7193,7 +7188,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i); decl = TREE_OPERAND (t, 0); var = create_tmp_var (TREE_TYPE (decl), get_name (decl)); - omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN); + omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN); TREE_OPERAND (t, 0) = var; t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i); TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1)); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index f8bacc3..a979d31 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,7 @@ +2014-11-12 Thomas Schwinge + + * c-c++-common/goacc/loop-private-1.c: New file. + 2014-11-07 Thomas Schwinge James Norris diff --git gcc/testsuite/c-c++-common/goacc/loop-private-1.c gcc/testsuite/c-c++-common/goacc/loop-private-1.c new file mode 100644 index 0000000..a54edb2 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/loop-private-1.c @@ -0,0 +1,14 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void +f (int i, int j) +{ +#pragma acc kernels +#pragma acc loop collapse(2) + for (i = 0; i < 20; ++i) + for (j = 0; j < 25; ++j) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc loop collapse\\(2\\) private\\(j\\) private\\(i\\)" 1 "gimple" } } */ +/* { dg-final { cleanup-tree-dump "gimple" } } */ diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 9f562d5..36beb59 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-11-12 Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file. + * testsuite/libgomp.oacc-c/collapse-4.c: Turn into an execution + test, check result. + * testsuite/libgomp.oacc-c-c++-common/data-2.c: Update. + * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. + 2014-11-07 Thomas Schwinge * plugin-nvptx.c: Remove. diff --git libgomp/testsuite/libgomp.oacc-c/collapse-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c similarity index 56% copy from libgomp/testsuite/libgomp.oacc-c/collapse-4.c copy to libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c index 7df0de2..f3538a6 100644 --- libgomp/testsuite/libgomp.oacc-c/collapse-4.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c @@ -1,6 +1,4 @@ -/* TODO: change to a run test once libgomp supports all data clauses. */ -/* { dg-do compile } */ -/* { dg-options "-O2 -std=c99" } */ +/* See also ../libgomp.oacc-c/collapse-4.c. */ #include @@ -9,17 +7,21 @@ main (void) { int l = 0; int b[3][3]; + int i, j; memset (b, '\0', sizeof (b)); #pragma acc parallel copy(b[0:3][0:3]) copy(l) { #pragma acc loop collapse(2) reduction(+:l) - for (int i = 0; i < 2; i++) - for (int j = 0; j < 2; j++) + for (i = 0; i < 2; i++) + for (j = 0; j < 2; j++) if (b[i][j] != 16) l += 1; } + if (l != 2 * 2) + __builtin_abort(); + return 0; } diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c index b990ade..3e02ffa 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c @@ -28,8 +28,8 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async #pragma acc parallel async wait #pragma acc loop - for (int ii = 0; ii < N; ii++) - b[ii] = a[ii]; + for (i = 0; i < N; i++) + b[i] = a[i]; #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async #pragma acc wait @@ -52,8 +52,8 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) #pragma acc parallel async (1) #pragma acc loop - for (int ii = 0; ii < N; ii++) - b[ii] = a[ii]; + for (i = 0; i < N; i++) + b[i] = a[i]; #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1) #pragma acc wait (1) @@ -79,18 +79,18 @@ main (int argc, char **argv) #pragma acc parallel async (1) wait (1) #pragma acc loop - for (int ii = 0; ii < N; ii++) - b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + for (i = 0; i < N; i++) + b[i] = (a[i] * a[i] * a[i]) / a[i]; #pragma acc parallel async (2) wait (1) #pragma acc loop - for (int ii = 0; ii < N; ii++) - c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + for (i = 0; i < N; i++) + c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; #pragma acc parallel async (3) wait (1) #pragma acc loop - for (int ii = 0; ii < N; ii++) - d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + for (i = 0; i < N; i++) + d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1) #pragma acc wait (1) diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c index f8f1b3b..44787dd 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c @@ -28,8 +28,8 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async #pragma acc parallel async wait #pragma acc loop - for (int ii = 0; ii < N; ii++) - b[ii] = a[ii]; + for (i = 0; i < N; i++) + b[i] = a[i]; #pragma acc update host (a[0:N], b[0:N]) async wait #pragma acc wait @@ -52,8 +52,8 @@ main (int argc, char **argv) #pragma acc update device (a[0:N], b[0:N]) async (1) #pragma acc parallel async (1) #pragma acc loop - for (int ii = 0; ii < N; ii++) - b[ii] = a[ii]; + for (i = 0; i < N; i++) + b[i] = a[i]; #pragma acc update host (a[0:N], b[0:N]) async (1) wait (1) #pragma acc wait (1) @@ -81,18 +81,18 @@ main (int argc, char **argv) #pragma acc parallel async (1) wait (1,2) #pragma acc loop - for (int ii = 0; ii < N; ii++) - b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + for (i = 0; i < N; i++) + b[i] = (a[i] * a[i] * a[i]) / a[i]; #pragma acc parallel async (2) wait (1,3) #pragma acc loop - for (int ii = 0; ii < N; ii++) - c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + for (i = 0; i < N; i++) + c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; #pragma acc parallel async (3) wait (1,3) #pragma acc loop - for (int ii = 0; ii < N; ii++) - d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + for (i = 0; i < N; i++) + d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; #pragma acc update host (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) wait (1,2,3) #pragma acc wait (1) diff --git libgomp/testsuite/libgomp.oacc-c/collapse-4.c libgomp/testsuite/libgomp.oacc-c/collapse-4.c index 7df0de2..17663ca 100644 --- libgomp/testsuite/libgomp.oacc-c/collapse-4.c +++ libgomp/testsuite/libgomp.oacc-c/collapse-4.c @@ -1,5 +1,4 @@ -/* TODO: change to a run test once libgomp supports all data clauses. */ -/* { dg-do compile } */ +/* See also ../libgomp.oacc-c-c++-common/collapse-4.c. */ /* { dg-options "-O2 -std=c99" } */ #include @@ -21,5 +20,8 @@ main (void) l += 1; } + if (l != 2 * 2) + __builtin_abort(); + return 0; }