From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 13865 invoked by alias); 26 Jul 2018 08:34:02 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 13843 invoked by uid 89); 26 Jul 2018 08:34:02 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-25.5 required=5.0 tests=AWL,BAYES_00,FREEMAIL_FROM,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy= X-HELO: mail-lf1-f66.google.com Received: from mail-lf1-f66.google.com (HELO mail-lf1-f66.google.com) (209.85.167.66) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 26 Jul 2018 08:33:59 +0000 Received: by mail-lf1-f66.google.com with SMTP id a4-v6so623962lff.5 for ; Thu, 26 Jul 2018 01:33:59 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20161025; h=mime-version:references:in-reply-to:from:date:message-id:subject:to :cc; bh=gGiY4nvhLpTUmV0SRwPy4Dy+33RWxbDx4tqWQBT1FT4=; b=hovpWSdA6gPWRmnsHNeqsWyO77Qil0DE0hvgwVeBySugy8B2oCOP2p03Jq6qp0a+pg /dkHDUo0gb8UbsBA0tnxgEfvnOQiw9Pp4P4exaFhKCOsVjiSL80CNuPirC3sydCtW1q+ wlC0+yYwUBIh+UtBACyK8vB5XqHo7JCFYIec3A9XGU+MAp75zqU3VNJVoACdZfObMq6S KEOKmDtjdsDfu7op+Yi4in8LlLYX6AIBJ8snOSPKR+cqsJk7mz8zCX5Wu9JuB3tAGh0k 89eKkhqU1CekBYrSaE17bzpukacib9MVcNYnPY3NUFIvUo+mNIO1YOOfS+D5i/eP0Ri5 UYgw== MIME-Version: 1.0 References: <71c14c108ca5d2c936c4350ed127aa2276316f1d.1532531520.git.cesar@codesourcery.com> In-Reply-To: <71c14c108ca5d2c936c4350ed127aa2276316f1d.1532531520.git.cesar@codesourcery.com> From: Richard Biener Date: Thu, 26 Jul 2018 08:34:00 -0000 Message-ID: Subject: Re: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism. To: cesar@codesourcery.com Cc: GCC Patches , Jakub Jelinek Content-Type: text/plain; charset="UTF-8" X-IsSubscribed: yes X-SW-Source: 2018-07/txt/msg01588.txt.bz2 On Wed, Jul 25, 2018 at 5:30 PM Cesar Philippidis wrote: > > This patch teaches GCC to inform the user how it assigned parallelism > to each OpenACC loop at compile time using the -fopt-info-note-omp > flag. For instance, given the acc parallel loop nest: > > #pragma acc parallel loop > for (...) > #pragma acc loop vector > for (...) > > GCC will report somthing like > > foo.c:4:0: note: Detected parallelism > foo.c:6:0: note: Detected parallelism > > Note how only the inner loop specifies vector parallelism. In this > example, GCC automatically assigned gang and worker parallelism to the > outermost loop. Perhaps, going forward, it would be useful to > distinguish which parallelism was specified by the user and which was > assigned by the compiler. But that can be added in a follow up patch. > > Is this patch OK for trunk? I bootstrapped and regtested it for x86_64 > with nvptx offloading. Shouldn't this use MSG_OPTIMIZED_LOCATIONS instead? Are there any other optinfo notes emitted? Like when despite pragmas loops are not handled or so? > Thanks, > Cesar > > 2018-XX-YY Cesar Philippidis > > gcc/ > * omp-offload.c (inform_oacc_loop): New function. > (execute_oacc_device_lower): Use it to display loop parallelism. > > gcc/testsuite/ > * c-c++-common/goacc/note-parallelism.c: New test. > * gfortran.dg/goacc/note-parallelism.f90: New test. > > (cherry picked from gomp-4_0-branch r245683, and gcc/testsuite/ parts of > r245770) > > diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c > index 0abf028..66b99bb 100644 > --- a/gcc/omp-offload.c > +++ b/gcc/omp-offload.c > @@ -866,6 +866,31 @@ debug_oacc_loop (oacc_loop *loop) > dump_oacc_loop (stderr, loop, 0); > } > > +/* Provide diagnostics on OpenACC loops LOOP, its siblings and its > + children. */ > + > +static void > +inform_oacc_loop (oacc_loop *loop) > +{ > + const char *seq = loop->mask == 0 ? " seq" : ""; > + const char *gang = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) > + ? " gang" : ""; > + const char *worker = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) > + ? " worker" : ""; > + const char *vector = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) > + ? " vector" : ""; > + dump_location_t loc = dump_location_t::from_location_t (loop->loc); > + > + dump_printf_loc (MSG_NOTE, loc, > + "Detected parallelism \n", seq, gang, > + worker, vector); > + > + if (loop->child) > + inform_oacc_loop (loop->child); > + if (loop->sibling) > + inform_oacc_loop (loop->sibling); > +} > + > /* DFS walk of basic blocks BB onwards, creating OpenACC loop > structures as we go. By construction these loops are properly > nested. */ > @@ -1533,6 +1558,8 @@ execute_oacc_device_lower () > dump_oacc_loop (dump_file, loops, 0); > fprintf (dump_file, "\n"); > } > + if (dump_enabled_p () && loops->child) > + inform_oacc_loop (loops->child); > > /* Offloaded targets may introduce new basic blocks, which require > dominance information to update SSA. */ > diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c > new file mode 100644 > index 0000000..3ec794c > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c > @@ -0,0 +1,61 @@ > +/* Test the output of -fopt-info-note-omp. */ > + > +/* { dg-additional-options "-fopt-info-note-omp" } */ > + > +int > +main () > +{ > + int x, y, z; > + > +#pragma acc parallel loop seq /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop worker /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop vector /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang vector /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang worker /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop worker vector /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang worker vector /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > +#pragma acc loop /* { dg-message "note: Detected parallelism " } */ > + for (y = 0; y < 10; y++) > + ; > + > +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism " } */ > + for (x = 0; x < 10; x++) > +#pragma acc loop worker /* { dg-message "note: Detected parallelism " } */ > + for (y = 0; y < 10; y++) > +#pragma acc loop vector /* { dg-message "note: Detected parallelism " } */ > + for (z = 0; z < 10; z++) > + ; > + > + return 0; > +} > diff --git a/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 > new file mode 100644 > index 0000000..a0c78c5 > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 > @@ -0,0 +1,62 @@ > +! Test the output of -fopt-info-note-omp. > + > +! { dg-additional-options "-fopt-info-note-omp" } > + > +program test > + implicit none > + > + integer x, y, z > + > + !$acc parallel loop seq ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop worker ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop vector ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang vector ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang worker ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop worker vector ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang worker vector ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + end do > + > + !$acc parallel loop ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + !$acc loop ! { dg-message "note: Detected parallelism " } > + do y = 1, 10 > + end do > + end do > + > + !$acc parallel loop gang ! { dg-message "note: Detected parallelism " } > + do x = 1, 10 > + !$acc loop worker ! { dg-message "note: Detected parallelism " } > + do y = 1, 10 > + !$acc loop vector ! { dg-message "note: Detected parallelism " } > + do z = 1, 10 > + end do > + end do > + end do > +end program test > -- > 2.7.4 >