From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 12654 invoked by alias); 28 Jul 2015 08:21:11 -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 12581 invoked by uid 89); 28 Jul 2015 08:21:10 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.7 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD,SPF_PASS autolearn=ham version=3.3.2 X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Tue, 28 Jul 2015 08:21:09 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:40846) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1ZK08I-0006YX-NP for gcc-patches@gnu.org; Tue, 28 Jul 2015 04:21:06 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1ZK08F-0002CM-1X for gcc-patches@gnu.org; Tue, 28 Jul 2015 04:21:06 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:45851) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1ZK08E-0002BV-S9 for gcc-patches@gnu.org; Tue, 28 Jul 2015 04:21:02 -0400 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZK08C-00064f-Uv from Tom_deVries@mentor.com ; Tue, 28 Jul 2015 01:21:01 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.224.2; Tue, 28 Jul 2015 09:20:59 +0100 Message-ID: <55B73B57.5090606@mentor.com> Date: Tue, 28 Jul 2015 09:10:00 -0000 From: Tom de Vries User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: "gcc-patches@gnu.org" , Jakub Jelinek Subject: [gomp4, committed] Handle double reduction in oacc kernels pass group Content-Type: multipart/mixed; boundary="------------020001030008090507010803" X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 X-SW-Source: 2015-07/txt/msg02331.txt.bz2 --------------020001030008090507010803 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 296 Hi, this patch adds a test-case with a double reduction in an oacc kernels region. In order to get it in the proper shape for parloops to deal with, I needed to repeat the pass_lim/pass_copy_prop sequence. Bootstrapped and reg-tested on x86_64. Committed to gomp-4_0-branch. Thanks, - Tom --------------020001030008090507010803 Content-Type: text/x-patch; name="0005-Handle-double-reduction-in-oacc-kernels-pass-group.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename*0="0005-Handle-double-reduction-in-oacc-kernels-pass-group.patc"; filename*1="h" Content-length: 2509 Handle double reduction in oacc kernels pass group 2015-07-28 Tom de Vries * passes.def: Repeat pass_lim and pass_copy_prop in oacc kernels pass group. * c-c++-common/goacc/kernels-double-reduction.c: New test. --- gcc/passes.def | 2 ++ .../c-c++-common/goacc/kernels-double-reduction.c | 37 ++++++++++++++++++++++ 2 files changed, 39 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c diff --git a/gcc/passes.def b/gcc/passes.def index ae91ed1..e31e39f 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -96,6 +96,8 @@ along with GCC; see the file COPYING3. If not see NEXT_PASS (pass_tree_loop_init); NEXT_PASS (pass_lim); NEXT_PASS (pass_copy_prop); + NEXT_PASS (pass_lim); + NEXT_PASS (pass_copy_prop); NEXT_PASS (pass_scev_cprop); NEXT_PASS (pass_parallelize_loops_oacc_kernels); NEXT_PASS (pass_expand_omp_ssa); diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c new file mode 100644 index 0000000..81467a9 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c @@ -0,0 +1,37 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ +/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */ +/* { dg-additional-options "-fdump-tree-optimized" } */ + +#include + +#define N 500 + +unsigned int a[N][N]; + +void __attribute__((noinline,noclone)) +foo (void) +{ + int i, j; + unsigned int sum = 1; + +#pragma acc kernels copyin (a[0:N]) copy (sum) + { + for (i = 0; i < N; ++i) + for (j = 0; j < N; ++j) + sum += a[i][j]; + } + + if (sum != 5001) + abort (); +} + +/* Check that only one loop is analyzed, and that it can be parallelized. */ +/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "parallelizing outer loop" 1 "parloops_oacc_kernels" } } */ + +/* Check that the loop has been split off into a function. */ +/* { dg-final { scan-tree-dump-times "(?n);; Function .*foo.*._omp_fn.0" 1 "optimized" } } */ + +/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */ -- 1.9.1 --------------020001030008090507010803--