From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 53234 invoked by alias); 12 Oct 2015 17:29:47 -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 53221 invoked by uid 89); 12 Oct 2015 17:29:47 -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,SPF_PASS,T_RP_MATCHES_RCVD 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; Mon, 12 Oct 2015 17:29:45 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:60512) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1Zlgus-0000kV-Uv for gcc-patches@gnu.org; Mon, 12 Oct 2015 13:29:43 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1Zlguo-0003Iu-9q for gcc-patches@gnu.org; Mon, 12 Oct 2015 13:29:42 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:55551) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1Zlguo-0003Io-1X for gcc-patches@gnu.org; Mon, 12 Oct 2015 13:29:38 -0400 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 1Zlgun-0006ei-6H from Tom_deVries@mentor.com ; Mon, 12 Oct 2015 10:29:37 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Mon, 12 Oct 2015 18:29:35 +0100 Subject: [committed, gomp4, 3/3] Handle sequential code in kernels region - Testcases To: "gcc-patches@gnu.org" References: <561BEA02.6010808@mentor.com> CC: Jakub Jelinek From: Tom de Vries Message-ID: <561BEDD1.7040203@mentor.com> Date: Mon, 12 Oct 2015 17:29:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <561BEA02.6010808@mentor.com> Content-Type: multipart/mixed; boundary="------------030400040700040605010605" X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 X-SW-Source: 2015-10/txt/msg01159.txt.bz2 --------------030400040700040605010605 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 618 On 12/10/15 19:12, Tom de Vries wrote: > Hi, > > I've committed the following patch series. > > 1 Add get_bbs_in_oacc_kernels_region > 2 Handle sequential code in kernels region > 3 Handle sequential code in kernels region - Testcases > > The patch series adds detection of whether sequential code (that is, > code in the oacc kernels region before and after the loop that is to be > parallelized), is safe to execute in parallel. > > Bootstrapped and reg-tested on x86_64. > > I'll post the patches individually, in reply to this email. This patch adds relevant test-cases. Thanks, - Tom --------------030400040700040605010605 Content-Type: text/x-patch; name="0003-Handle-sequential-code-in-kernels-region-Testcases.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename*0="0003-Handle-sequential-code-in-kernels-region-Testcases.patc"; filename*1="h" Content-length: 6415 Handle sequential code in kernels region - Testcases 2015-10-12 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c: New test. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c: New test. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c: New test. --- .../kernels-loop-and-seq-2.c | 36 +++++++++++++++++++++ .../kernels-loop-and-seq-3.c | 37 ++++++++++++++++++++++ .../kernels-loop-and-seq-4.c | 36 +++++++++++++++++++++ .../kernels-loop-and-seq-5.c | 37 ++++++++++++++++++++++ .../kernels-loop-and-seq-6.c | 36 +++++++++++++++++++++ .../kernels-loop-and-seq.c | 37 ++++++++++++++++++++++ 6 files changed, 219 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c new file mode 100644 index 0000000..2e4100f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c @@ -0,0 +1,36 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 32 + +unsigned int +foo (int n, unsigned int *a) +{ +#pragma acc kernels copy (a[0:N]) + { + a[0] = a[0] + 1; + + for (int i = 0; i < n; i++) + a[i] = 1; + } + + return a[0]; +} + +int +main (void) +{ + unsigned int a[N]; + unsigned res, i; + + for (i = 0; i < N; ++i) + a[i] = i % 4; + + res = foo (N, a); + if (res != 1) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c new file mode 100644 index 0000000..b3e736b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 32 + +unsigned int +foo (int n, unsigned int *a) +{ + +#pragma acc kernels copy (a[0:N]) + { + for (int i = 0; i < n; i++) + a[i] = 1; + + a[0] = 2; + } + + return a[0]; +} + +int +main (void) +{ + unsigned int a[N]; + unsigned res, i; + + for (i = 0; i < N; ++i) + a[i] = i % 4; + + res = foo (N, a); + if (res != 2) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c new file mode 100644 index 0000000..8b9affa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c @@ -0,0 +1,36 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 32 + +unsigned int +foo (int n, unsigned int *a) +{ +#pragma acc kernels copy (a[0:N]) + { + a[0] = 2; + + for (int i = 0; i < n; i++) + a[i] = 1; + } + + return a[0]; +} + +int +main (void) +{ + unsigned int a[N]; + unsigned res, i; + + for (i = 0; i < N; ++i) + a[i] = i % 4; + + res = foo (N, a); + if (res != 1) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c new file mode 100644 index 0000000..83d4e7f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 32 + +unsigned int +foo (int n, unsigned int *a) +{ + int r; +#pragma acc kernels copyout(r) copy (a[0:N]) + { + r = a[0]; + + for (int i = 0; i < n; i++) + a[i] = 1; + } + + return r; +} + +int +main (void) +{ + unsigned int a[N]; + unsigned res, i; + + for (i = 0; i < N; ++i) + a[i] = i % 4; + + res = foo (N, a); + if (res != 0) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c new file mode 100644 index 0000000..01d5e5e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c @@ -0,0 +1,36 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 32 + +unsigned int +foo (int n, unsigned int *a) +{ +#pragma acc kernels copy (a[0:N]) + { + int r = a[0]; + + for (int i = 0; i < n; i++) + a[i] = 1 + r; + } + + return a[0]; +} + +int +main (void) +{ + unsigned int a[N]; + unsigned res, i; + + for (i = 0; i < N; ++i) + a[i] = i % 4; + + res = foo (N, a); + if (res != 1) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c new file mode 100644 index 0000000..61d1283 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 32 + +unsigned int +foo (int n, unsigned int *a) +{ + +#pragma acc kernels copy (a[0:N]) + { + for (int i = 0; i < n; i++) + a[i] = 1; + + a[0] = a[0] + 1; + } + + return a[0]; +} + +int +main (void) +{ + unsigned int a[N]; + unsigned res, i; + + for (i = 0; i < N; ++i) + a[i] = i % 4; + + res = foo (N, a); + if (res != 2) + abort (); + + return 0; +} -- 1.9.1 --------------030400040700040605010605--