From patchwork Mon Jan 18 13:39:05 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 569545 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 E92EE1402D6 for ; Tue, 19 Jan 2016 00:39:36 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=icNJjyAb; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=FTj4rE3HFKCNtY6A2 ckMW5zMSOLG8ZZaYGwnMeehBnQA5wWijGYFOTlYLVbJD9o5z7MRxIdK1Pag3wj8b 5I52xrwQX7H+qYhgqr3g3LixFcpUU9cPGu5HLd3rCiMhnq79gGvEYonYWwueUbUz 80InS5wEsDweqsjnuQuWfed9oQ= 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=R19dfrly0Sjdj+LIG6FtIa8 Cod0=; b=icNJjyAbyiiwuMQy5Ii7KfzliHNxnkBm1gagOUqnGTFrbOkWOXIpu0U KlP6MP8PVj1OCU3ESIUDtSyym6uyH5IjMj/UEFMILMH5BzlR8SYNQUOQ5Tac9+qn tV8uJIrrfi+A0LJD2mKUX29wHNzUU+RycCuM88vDvIRp/V7EjerI= Received: (qmail 57902 invoked by alias); 18 Jan 2016 13:39:28 -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 57888 invoked by uid 89); 18 Jan 2016 13:39:27 -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, RP_MATCHES_RCVD, SPF_PASS autolearn=ham version=3.3.2 spammy=1, 34, 1, 40, 537, HX-detected-operating-system:Windows 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, 18 Jan 2016 13:39:24 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:45795) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1aLA1i-0006kG-2I for gcc-patches@gnu.org; Mon, 18 Jan 2016 08:39:22 -0500 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1aLA1c-00037m-OQ for gcc-patches@gnu.org; Mon, 18 Jan 2016 08:39:21 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:55858) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1aLA1c-00037Z-CJ for gcc-patches@gnu.org; Mon, 18 Jan 2016 08:39:16 -0500 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 1aLA1a-0003ap-MN from Tom_deVries@mentor.com ; Mon, 18 Jan 2016 05:39:15 -0800 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; Mon, 18 Jan 2016 13:39:13 +0000 Subject: [comitted] Add oacc kernels test in libgomp To: "gcc-patches@gnu.org" References: <5640BD31.2060602@mentor.com> <5640FDB6.70903@mentor.com> CC: Jakub Jelinek , Richard Biener From: Tom de Vries Message-ID: <569CEAF9.1020003@mentor.com> Date: Mon, 18 Jan 2016 14:39:05 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.5.1 MIME-Version: 1.0 In-Reply-To: <5640FDB6.70903@mentor.com> X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 [ was: Re: [PATCH, 15/16] Add libgomp.oacc-c-c++-common/kernels-*.c ] On 09/11/15 21:10, Tom de Vries wrote: > On 09/11/15 16:35, Tom de Vries wrote: >> Hi, >> >> this patch series for stage1 trunk adds support to: >> - parallelize oacc kernels regions using parloops, and >> - map the loops onto the oacc gang dimension. >> >> The patch series contains these patches: >> >> 1 Insert new exit block only when needed in >> transform_to_exit_first_loop_alt >> 2 Make create_parallel_loop return void >> 3 Ignore reduction clause on kernels directive >> 4 Implement -foffload-alias >> 5 Add in_oacc_kernels_region in struct loop >> 6 Add pass_oacc_kernels >> 7 Add pass_dominator_oacc_kernels >> 8 Add pass_ch_oacc_kernels >> 9 Add pass_parallelize_loops_oacc_kernels >> 10 Add pass_oacc_kernels pass group in passes.def >> 11 Update testcases after adding kernels pass group >> 12 Handle acc loop directive >> 13 Add c-c++-common/goacc/kernels-*.c >> 14 Add gfortran.dg/goacc/kernels-*.f95 >> 15 Add libgomp.oacc-c-c++-common/kernels-*.c >> 16 Add libgomp.oacc-fortran/kernels-*.f95 >> >> The first 9 patches are more or less independent, but patches 10-16 are >> intended to be committed at the same time. >> >> Bootstrapped and reg-tested on x86_64. >> >> Build and reg-tested with nvidia accelerator, in combination with a >> patch that enables accelerator testing (which is submitted at >> https://gcc.gnu.org/ml/gcc-patches/2015-10/msg01771.html ). >> >> I'll post the individual patches in reply to this message. > > This patch adds C/C++ oacc kernels execution tests. > Bootstrapped and reg-tested on x86_64. Build with nvidia accelerator and tested goacc.exp and libgomp. Committed to trunk as attached (AFAICT, no changes compared to original posting, other than commit title). Thanks, - Tom Add oacc kernels test in libgomp 2015-11-09 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-2.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit-2.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-update.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-data.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-loop.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c: Same. * testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c: Same. --- .../libgomp.oacc-c-c++-common/kernels-loop-2.c | 47 ++++++++++++++++++++++ .../libgomp.oacc-c-c++-common/kernels-loop-3.c | 34 ++++++++++++++++ .../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 +++++++++++++++++ .../kernels-loop-collapse.c | 40 ++++++++++++++++++ .../libgomp.oacc-c-c++-common/kernels-loop-g.c | 5 +++ .../kernels-loop-mod-not-zero.c | 41 +++++++++++++++++++ .../libgomp.oacc-c-c++-common/kernels-loop-n.c | 47 ++++++++++++++++++++++ .../libgomp.oacc-c-c++-common/kernels-loop-nest.c | 26 ++++++++++++ .../libgomp.oacc-c-c++-common/kernels-loop.c | 41 +++++++++++++++++++ .../libgomp.oacc-c-c++-common/kernels-reduction.c | 37 +++++++++++++++++ 15 files changed, 537 insertions(+) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c new file mode 100644 index 0000000..13e57bd --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N (1024 * 512) +#define COUNTERTYPE unsigned int + +int +main (void) +{ + unsigned int *__restrict a; + unsigned int *__restrict b; + unsigned int *__restrict c; + + a = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + +#pragma acc kernels copyout (a[0:N]) + { + for (COUNTERTYPE i = 0; i < N; i++) + a[i] = i * 2; + } + +#pragma acc kernels copyout (b[0:N]) + { + for (COUNTERTYPE i = 0; i < N; i++) + b[i] = i * 4; + } + +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) + { + for (COUNTERTYPE ii = 0; ii < N; ii++) + c[ii] = a[ii] + b[ii]; + } + + for (COUNTERTYPE i = 0; i < N; i++) + if (c[i] != a[i] + b[i]) + abort (); + + free (a); + free (b); + free (c); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c new file mode 100644 index 0000000..f61a74a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c @@ -0,0 +1,34 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N (1024 * 512) +#define COUNTERTYPE unsigned int + +int +main (void) +{ + unsigned int i; + + unsigned int *__restrict c; + + c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + + for (COUNTERTYPE i = 0; i < N; i++) + c[i] = i * 2; + +#pragma acc kernels copy (c[0:N]) + { + for (COUNTERTYPE ii = 0; ii < N; ii++) + c[ii] = c[ii] + ii + 1; + } + + for (COUNTERTYPE i = 0; i < N; i++) + if (c[i] != i * 2 + i + 1) + abort (); + + free (c); + + return 0; +} 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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c new file mode 100644 index 0000000..f7f04cb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c @@ -0,0 +1,40 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 100 + +int a[N][N]; + +void __attribute__((noinline, noclone)) +foo (int m, int n) +{ + int i, j; + #pragma acc kernels + { +#pragma acc loop collapse(2) + for (i = 0; i < m; i++) + for (j = 0; j < n; j++) + a[i][j] = 1; + } +} + +int +main (void) +{ + int i, j; + + for (i = 0; i < N; i++) + for (j = 0; j < N; j++) + a[i][j] = 0; + + foo (N, N); + + for (i = 0; i < N; i++) + for (j = 0; j < N; j++) + if (a[i][j] != 1) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c new file mode 100644 index 0000000..96b6e4e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c @@ -0,0 +1,5 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ +/* { dg-additional-options "-g" } */ + +#include "kernels-loop.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c new file mode 100644 index 0000000..1433cb2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N ((1024 * 512) + 1) +#define COUNTERTYPE unsigned int + +int +main (void) +{ + unsigned int *__restrict a; + unsigned int *__restrict b; + unsigned int *__restrict c; + + a = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + + for (COUNTERTYPE i = 0; i < N; i++) + a[i] = i * 2; + + for (COUNTERTYPE i = 0; i < N; i++) + b[i] = i * 4; + +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) + { + for (COUNTERTYPE ii = 0; ii < N; ii++) + c[ii] = a[ii] + b[ii]; + } + + for (COUNTERTYPE i = 0; i < N; i++) + if (c[i] != a[i] + b[i]) + abort (); + + free (a); + free (b); + free (c); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c new file mode 100644 index 0000000..fd0d5b1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N ((1024 * 512) + 1) +#define COUNTERTYPE unsigned int + +static int __attribute__((noinline,noclone)) +foo (COUNTERTYPE n) +{ + unsigned int *__restrict a; + unsigned int *__restrict b; + unsigned int *__restrict c; + + a = (unsigned int *__restrict)malloc (n * sizeof (unsigned int)); + b = (unsigned int *__restrict)malloc (n * sizeof (unsigned int)); + c = (unsigned int *__restrict)malloc (n * sizeof (unsigned int)); + + for (COUNTERTYPE i = 0; i < n; i++) + a[i] = i * 2; + + for (COUNTERTYPE i = 0; i < n; i++) + b[i] = i * 4; + +#pragma acc kernels copyin (a[0:n], b[0:n]) copyout (c[0:n]) + { + for (COUNTERTYPE ii = 0; ii < n; ii++) + c[ii] = a[ii] + b[ii]; + } + + for (COUNTERTYPE i = 0; i < n; i++) + if (c[i] != a[i] + b[i]) + abort (); + + free (a); + free (b); + free (c); + + return 0; +} + +int +main (void) +{ + return foo (N); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c new file mode 100644 index 0000000..21d2599 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c @@ -0,0 +1,26 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N 1000 + +int +main (void) +{ + int x[N][N]; + +#pragma acc kernels copyout (x) + { + for (int ii = 0; ii < N; ii++) + for (int jj = 0; jj < N; jj++) + x[ii][jj] = ii + jj + 3; + } + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + if (x[i][j] != i + j + 3) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop.c new file mode 100644 index 0000000..3762e5a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define N (1024 * 512) +#define COUNTERTYPE unsigned int + +int +main (void) +{ + unsigned int *__restrict a; + unsigned int *__restrict b; + unsigned int *__restrict c; + + a = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); + + for (COUNTERTYPE i = 0; i < N; i++) + a[i] = i * 2; + + for (COUNTERTYPE i = 0; i < N; i++) + b[i] = i * 4; + +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) + { + for (COUNTERTYPE ii = 0; ii < N; ii++) + c[ii] = a[ii] + b[ii]; + } + + for (COUNTERTYPE i = 0; i < N; i++) + if (c[i] != a[i] + b[i]) + abort (); + + free (a); + free (b); + free (c); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c new file mode 100644 index 0000000..511e25f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#define n 10000 + +unsigned int a[n]; + +void __attribute__((noinline,noclone)) +foo (void) +{ + int i; + unsigned int sum = 1; + +#pragma acc kernels copyin (a[0:n]) copy (sum) + { + for (i = 0; i < n; ++i) + sum += a[i]; + } + + if (sum != 5001) + abort (); +} + +int +main () +{ + int i; + + for (i = 0; i < n; ++i) + a[i] = i % 2; + + foo (); + + return 0; +}