From patchwork Wed Jul 17 21:32:06 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 1133454 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-505236-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="dQZN7HUw"; dkim-atps=neutral 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 45pr8T5flBz9sDQ for ; Thu, 18 Jul 2019 07:32:41 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:from:to:cc:references:message-id:date:mime-version :in-reply-to:content-type:content-transfer-encoding; q=dns; s= default; b=NjWu7YICy9YF32NlOojuwZD3D6+jvFZtRiUW1iDf00m6vuXLk44gR 1LN5xuA8xyidIvtGTrQm43gGMtK0/V3VJZvQiC95MuSAzVxjGMzgcd27KhCI4erY BKMQe09tbsY4i4wy8Rc2R4C3lMB1yxStTm5FV/npEEGV/J8BrVWiyA= 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:from:to:cc:references:message-id:date:mime-version :in-reply-to:content-type:content-transfer-encoding; s=default; bh=3TzM6TVJq0Z912dP48+h94fQ3ok=; b=dQZN7HUwXiknCPkksKGK5BLHi4qv GTjdE1rVEULVdDi/yL7ZYIhp1rUoQyh9HOHRQBwSaTI+9csiH+PzguurNXK+KO0V Z0gTip/BRjTG4xN1jZMwnlacpF4XOW5HDj8X4v84rHchMDVF88V8jvG6vjzvY0ty rDVGG5kXzkdWXbM= Received: (qmail 43402 invoked by alias); 17 Jul 2019 21:32:29 -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 43212 invoked by uid 89); 17 Jul 2019 21:32:29 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-18.8 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.1 spammy= 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, 17 Jul 2019 21:32:18 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1hnrX5-0002p1-Qc from Kwok_Yeung@mentor.com ; Wed, 17 Jul 2019 14:32:16 -0700 Received: from [172.30.64.32] (137.202.0.90) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Wed, 17 Jul 2019 22:32:11 +0100 Subject: [PATCH 10/10, OpenACC] Make new OpenACC kernels conversion the default; adjust and add tests From: Kwok Cheung Yeung To: , Jakub Jelinek CC: Thomas Schwinge References: <5e191259-d5d5-34ce-7fd5-fc8d2e6d982e@codesourcery.com> Message-ID: Date: Wed, 17 Jul 2019 22:32:06 +0100 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:60.0) Gecko/20100101 Thunderbird/60.8.0 MIME-Version: 1.0 In-Reply-To: <5e191259-d5d5-34ce-7fd5-fc8d2e6d982e@codesourcery.com> This patch makes the new kernel conversion scheme the default, and adjusts the tests accordingly. 2019-07-16 Thomas Schwinge Kwok Cheung Yeung gcc/c-family/ * c.opt (fopenacc-kernels): Default to "split". gcc/fortran/ * lang.opt (fopenacc-kernels): Default to "split". gcc/ * doc/invoke.texi (-fopenacc-kernels): Update. gcc/testsuite/ * c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c: New file. * c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c: Likewise. * c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-1-kernels-loops.c: Likewise. * c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c: Likewise. * c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c: Likewise. * c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-loop-auto.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-loops.c: Likewise. * c-c++-common/goacc/classify-kernels-unparallelized.c: Update. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/classify-parallel.c: Likewise. * c-c++-common/goacc/classify-routine.c: Likewise. * c-c++-common/goacc/if-clause-2.c: Likewise. * c-c++-common/goacc/kernels-conversion.c: Likewise. * c-c++-common/goacc/kernels-decompose-1.c: Likewise. * c-c++-common/goacc/loop-2-kernels.c: Likewise. * c-c++-common/goacc/note-parallelism.c: Likewise. * c-c++-common/goacc/uninit-dim-clause.c: Likewise. * gfortran.dg/goacc/kernels-conversion.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/classify-kernels-unparallelized.f95 * gfortran.dg/goacc/classify-kernels.f95 * gfortran.dg/goacc/loop-2-kernels.f95 libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. --- gcc/c-family/c.opt | 2 +- gcc/doc/invoke.texi | 2 +- gcc/fortran/lang.opt | 2 +- .../goacc/classify-kernels-unparallelized.c | 9 +- .../c-c++-common/goacc/classify-kernels.c | 4 +- .../c-c++-common/goacc/classify-parallel.c | 2 +- .../c-c++-common/goacc/classify-routine.c | 2 +- gcc/testsuite/c-c++-common/goacc/if-clause-2.c | 1 - .../c-c++-common/goacc/kernels-conversion.c | 10 +- .../c-c++-common/goacc/kernels-decompose-1.c | 69 ++++--- gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c | 14 +- ...sm-1-kernels-conditional-loop-independent_seq.c | 129 +++++++++++++ .../goacc/note-parallelism-1-kernels-loop-auto.c | 126 +++++++++++++ ...te-parallelism-1-kernels-loop-independent_seq.c | 126 +++++++++++++ .../goacc/note-parallelism-1-kernels-loops.c | 47 +++++ .../note-parallelism-1-kernels-straight-line.c | 82 +++++++++ .../note-parallelism-combined-kernels-loop-auto.c | 121 ++++++++++++ ...llelism-combined-kernels-loop-independent_seq.c | 121 ++++++++++++ ...lism-kernels-conditional-loop-independent_seq.c | 204 +++++++++++++++++++++ .../goacc/note-parallelism-kernels-loop-auto.c | 138 ++++++++++++++ ...note-parallelism-kernels-loop-independent_seq.c | 138 ++++++++++++++ .../goacc/note-parallelism-kernels-loops.c | 50 +++++ .../c-c++-common/goacc/note-parallelism.c | 3 +- .../c-c++-common/goacc/uninit-dim-clause.c | 6 +- .../goacc/classify-kernels-unparallelized.f95 | 1 + .../gfortran.dg/goacc/classify-kernels.f95 | 1 + .../gfortran.dg/goacc/kernels-conversion.f95 | 7 +- .../gfortran.dg/goacc/kernels-decompose-1.f95 | 79 ++++---- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 1 - gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95 | 22 +-- .../libgomp.oacc-c-c++-common/acc_prof-kernels-1.c | 17 +- .../kernels-decompose-1.c | 9 +- 32 files changed, 1416 insertions(+), 129 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c b[i] = c; - a = c; /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + a = c; /* { dg-warning "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ } for (int i = 0; i < N; ++i) diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt index a193875..8efc5ea 100644 --- a/gcc/c-family/c.opt +++ b/gcc/c-family/c.opt @@ -1689,7 +1689,7 @@ C ObjC C++ ObjC++ LTO Joined Var(flag_openacc_dims) Specify default OpenACC compute dimensions. fopenacc-kernels= -C ObjC C++ ObjC++ RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_PARLOOPS) +C ObjC C++ ObjC++ RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_SPLIT) -fopenacc-kernels=[split|parloops] Configure OpenACC 'kernels' constructs handling. Enum diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index ec98ab6..ffde9a2 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -2200,9 +2200,9 @@ Configure OpenACC 'kernels' constructs handling. With @option{-fopenacc-kernels=split}, OpenACC 'kernels' constructs are split into a sequence of compute constructs, each then handled individually. +This is the default. With @option{-fopenacc-kernels=parloops}, the whole OpenACC 'kernels' constructs is handled by the @samp{parloops} pass. -This is the default. @item -fopenmp @opindex fopenmp diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt index e7e277a..c84b284 100644 --- a/gcc/fortran/lang.opt +++ b/gcc/fortran/lang.opt @@ -663,7 +663,7 @@ Fortran LTO Joined Var(flag_openacc_dims) ; Documented in C fopenacc-kernels= -Fortran RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_PARLOOPS) +Fortran RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_SPLIT) ; Documented in C fopenmp diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index d4c4b2c..443b207 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for unparallelized - OpenACC kernels. */ + OpenACC 'kernels'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } @@ -13,14 +13,15 @@ extern unsigned int *__restrict a; extern unsigned int *__restrict b; extern unsigned int *__restrict c; -/* An "extern"al mapping of loop iterations/array indices makes the loop - unparallelizable. */ extern unsigned int f (unsigned int); +#pragma acc routine (f) seq void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ - for (unsigned int i = 0; i < N; i++) + for (unsigned int i = 0; i < N; i++) /* { dg-message "optimized: beginning \"parloops\" region in OpenACC 'kernels' construct" } */ + /* An "extern"al mapping of loop iterations/array indices makes the loop + unparallelizable. */ c[i] = a[f (i)] + b[f (i)]; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 16e9b9e..c154edf 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for OpenACC - kernels. */ + 'kernels'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } @@ -16,7 +16,7 @@ extern unsigned int *__restrict c; void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ - for (unsigned int i = 0; i < N; i++) + for (unsigned int i = 0; i < N; i++) /* { dg-message "optimized: beginning \"parloops\" region in OpenACC 'kernels' construct" } */ c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 66a6d13..9c80efd 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for OpenACC - parallel. */ + 'parallel'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index 0b9ba6e..a4994b0 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for OpenACC - routine. */ + 'routine'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c index e17b5dd..9920b4f 100644 --- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c @@ -1,4 +1,3 @@ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */ void diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c b/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c index ea7eec9..8cb63f0 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c @@ -1,4 +1,3 @@ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */ #define N 1024 @@ -52,12 +51,11 @@ main (void) parallelized loop region; and three "old-style" kernel regions. */ /* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single" 1 "convert_oacc_kernels" } } */ /* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized" 1 "convert_oacc_kernels" } } */ -/* { dg-final { scan-tree-dump-times "oacc_kernels" 3 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_kernels " 3 "convert_oacc_kernels" } } */ /* Each of the parallel regions is async, and there is a final call to __builtin_GOACC_wait. */ -/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels.* async\(-1\)" 5 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\)" 3 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\)" 1 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized async\\(-1\\)" 1 "convert_oacc_kernels" } } */ /* { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1 "convert_oacc_kernels" } } */ - -/* Check that the original kernels region is removed. */ -/* { dg-final { scan-tree-dump-not "oacc_kernels" "convert_oacc_kernels" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c index b5d58c3..293ee42 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c @@ -1,6 +1,5 @@ /* Test OpenACC 'kernels' construct decomposition. */ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-O2" } for "parloops". */ @@ -31,92 +30,92 @@ main () #pragma acc kernels { - x = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + x = 0; /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ y = x < 10; z = x++; ; } -#pragma acc kernels - for (int i = 0; i < N; i++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + for (int i = 0; i < N; i++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ a[i] = 0; -#pragma acc kernels loop - /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ +#pragma acc kernels loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ for (int i = 0; i < N; i++) b[i] = a[N - i - 1]; #pragma acc kernels { -#pragma acc loop - /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ for (int i = 0; i < N; i++) b[i] = a[N - i - 1]; -#pragma acc loop - /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ for (int i = 0; i < N; i++) c[i] = a[i] * b[i]; - a[z] = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + a[z] = 0; /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ -#pragma acc loop - /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ for (int i = 0; i < N; i++) c[i] += a[i]; -#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ - /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ for (int i = 0 + 1; i < N; i++) c[i] += c[i - 1]; } -#pragma acc kernels +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ { -#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ - /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ for (int i = 0; i < N; ++i) -#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ for (int j = 0; j < N; ++j) -#pragma acc loop independent /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */ for (int k = 0; k < N; ++k) a[(i + j + k) % N] = b[j] - + f_v (c[k]); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + + f_v (c[k]); /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ //TODO Should the following turn into "gang-single" instead of "parloops"? //TODO The problem is that the first STMT is "if (y <= 4) goto ; else goto ;", thus "parloops". - if (y < 5) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ -#pragma acc loop independent /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" } */ + if (y < 5) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ +#pragma acc loop independent /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" } */ for (int j = 0; j < N; ++j) b[j] = f_w (c[j]); } -#pragma acc kernels /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */ +#pragma acc kernels { - /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" "" { target *-*-* } .+1 } */ - y = f_g (a[5]); /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" "" { target *-*-* } .+1 } */ + y = f_g (a[5]); /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ -#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ - /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ for (int j = 0; j < N; ++j) - b[j] = y + f_w (c[j]); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + b[j] = y + f_w (c[j]); /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ } #pragma acc kernels { - y = 3; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + y = 3; /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ -#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ - /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ for (int j = 0; j < N; ++j) - b[j] = y + f_v (c[j]); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + b[j] = y + f_v (c[j]); /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ - z = 2; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + z = 2; /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ } -#pragma acc kernels /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ +#pragma acc kernels /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ ; return 0; diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c index 0151508..c989222 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c @@ -37,7 +37,7 @@ void K(void) for (j = 0; j < 10; j++) { } } -#pragma acc loop seq gang // { dg-error "'seq' overrides" } +#pragma acc loop seq gang // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } @@ -63,7 +63,7 @@ void K(void) for (j = 0; j < 10; j++) { } } -#pragma acc loop seq worker // { dg-error "'seq' overrides" } +#pragma acc loop seq worker // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc loop gang worker @@ -92,7 +92,7 @@ void K(void) for (j = 1; j < 10; j++) { } } -#pragma acc loop seq vector // { dg-error "'seq' overrides" } +#pragma acc loop seq vector // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc loop gang vector @@ -105,7 +105,7 @@ void K(void) #pragma acc loop auto for (i = 0; i < 10; i++) { } -#pragma acc loop seq auto // { dg-error "'seq' overrides" } +#pragma acc loop seq auto // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc loop gang auto // { dg-error "'auto' conflicts" } @@ -147,7 +147,7 @@ void K(void) #pragma acc kernels loop worker(num:5) for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" } +#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang worker @@ -163,7 +163,7 @@ void K(void) #pragma acc kernels loop vector(length:5) for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" } +#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang vector @@ -176,7 +176,7 @@ void K(void) #pragma acc kernels loop auto for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" } +#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" } diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c new file mode 100644 index 0000000..c21273a --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c @@ -0,0 +1,129 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing conditionally executed 'loop' constructs with + 'independent' or 'seq' clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +extern int c; + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker + for (y = 0; y < 10; y++) +#pragma acc loop independent vector + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) + ; + +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c new file mode 100644 index 0000000..eedc472 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c @@ -0,0 +1,126 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing 'loop' constructs with explicit or implicit 'auto' + clause. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels + /* Strangely indented to keep this similar to other test cases. */ + { +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop gang /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop gang vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop gang worker /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop worker vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop gang worker vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop gang /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker + for (y = 0; y < 10; y++) +#pragma acc loop vector + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) + ; + +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c new file mode 100644 index 0000000..dad1bdb --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c @@ -0,0 +1,126 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing 'loop' constructs with 'independent' or 'seq' + clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels + /* Strangely indented to keep this similar to other test cases. */ + { +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang vector /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker vector /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker vector /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c new file mode 100644 index 0000000..336be88 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c @@ -0,0 +1,47 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing loops. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + { + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + + for (x = 0; x < 10; x++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + + for (x = 0; x < 10; x++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c new file mode 100644 index 0000000..07a1e32 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c @@ -0,0 +1,82 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing straight-line code. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +#pragma acc routine gang +extern int +f_g (int); + +#pragma acc routine worker +extern int +f_w (int); + +#pragma acc routine vector +extern int +f_v (int); + +#pragma acc routine seq +extern int +f_s (int); + +int +main () +{ + int x, y, z; + +#pragma acc kernels + { + x = 0; /* { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ + y = x < 10; + z = x++; + ; + + y = 0; + z = y < 10; + x -= f_g (y++); /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + ; + + x = f_w (0); /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + z = f_v (x < 10); /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + y -= f_s (x++); /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + ; + + x = 0; + y = x < 10; + z = (x++); + y = 0; + x = y < 10; + z += (y++); + ; + + x = 0; + y += f_s (x < 10); /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + x++; + y = 0; + y += f_v (y < 10); /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + y++; + z = 0; + y += f_w (z < 10); /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + z++; + ; + + x = 0; + y *= f_g ( /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + f_w (x < 10) /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + + f_g (x < 10) /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + ); + x++; + y = 0; + y *= y < 10; + y++; + z = 0; + y *= z < 10; + z++; + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c new file mode 100644 index 0000000..2241901 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c @@ -0,0 +1,121 @@ +/* Test the output of "-fopt-info-optimized-omp" for combined OpenACC 'kernels + loop' constructs with explicit or implicit 'auto' clause. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop gang /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop worker /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop gang vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop gang worker /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop worker vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop gang worker vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop gang /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker + for (y = 0; y < 10; y++) +#pragma acc loop vector + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c new file mode 100644 index 0000000..b743636 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c @@ -0,0 +1,121 @@ +/* Test the output of "-fopt-info-optimized-omp" for combined OpenACC 'kernels + loop' constructs with 'independent' or 'seq' clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang vector /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang worker /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent worker vector /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang worker vector /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels loop independent /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c new file mode 100644 index 0000000..21d31c4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c @@ -0,0 +1,204 @@ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels' + constructs containing conditionally executed 'loop' constructs with + 'independent' or 'seq' clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +extern int c; + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent worker + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang worker + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent worker vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang worker vector + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker + for (y = 0; y < 10; y++) +#pragma acc loop independent vector + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c new file mode 100644 index 0000000..02b9064 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c @@ -0,0 +1,138 @@ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels' + constructs containing 'loop' constructs with explicit or implicit 'auto' + clause. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop gang /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop gang vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop gang worker /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop worker vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop gang worker vector /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop gang /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker + for (y = 0; y < 10; y++) +#pragma acc loop vector + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c new file mode 100644 index 0000000..6824d70 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c @@ -0,0 +1,138 @@ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels' + constructs containing 'loop' constructs with 'independent' or 'seq' + clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang vector /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang worker /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent worker vector /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang worker vector /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c new file mode 100644 index 0000000..365464b --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c @@ -0,0 +1,50 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing loops. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c index 735df7d..2b49a8b 100644 --- a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c @@ -1,4 +1,5 @@ -/* Test the output of "-fopt-info-optimized-omp". */ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'parallel' + constructs. */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c index 9f11196..f00daa7 100644 --- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c +++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c @@ -18,12 +18,12 @@ void acc_kernels() { int i, j, k; - #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */ + #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */ ; - #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */ + #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */ ; - #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */ + #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */ ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index 0877242..27ba39b 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -21,6 +21,7 @@ program main !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-message "optimized: beginning \"parloops\" region in OpenACC 'kernels' construct" "" { target *-*-* } .-1 } c(i) = a(f (i)) + b(f (i)) end do !$acc end kernels diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index f2c4736..68d0512 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -17,6 +17,7 @@ program main !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + ! { dg-message "optimized: beginning \"parloops\" region in OpenACC 'kernels' construct" "" { target *-*-* } .-1 } c(i) = a(i) + b(i) end do !$acc end kernels diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 index 6604727..4672d15 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 @@ -1,4 +1,3 @@ -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fdump-tree-convert_oacc_kernels" } program main @@ -50,9 +49,11 @@ end program main ! parallelized loop region; and three "old-style" kernel regions. ! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single" 1 "convert_oacc_kernels" } } ! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized" 1 "convert_oacc_kernels" } } -! { dg-final { scan-tree-dump-times "oacc_kernels" 3 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_kernels " 3 "convert_oacc_kernels" } } ! Each of the parallel regions is async, and there is a final call to ! __builtin_GOACC_wait. -! { dg-final { scan-tree-dump-times "oacc_parallel_kernels.* async\(-1\)" 5 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\)" 3 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\)" 1 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized async\\(-1\\)" 1 "convert_oacc_kernels" } } ! { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1 "convert_oacc_kernels" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 index 520bf03..b2956d7 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 @@ -1,6 +1,5 @@ ! Test OpenACC 'kernels' construct decomposition. -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-O2" } for "parloops". @@ -25,7 +24,7 @@ program main integer :: a(N), b(N), c(N) !$acc kernels - x = 0 ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } + x = 0 ! { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } y = 0 y_l = x < 10 z = x @@ -33,67 +32,67 @@ program main ; !$acc end kernels - !$acc kernels ! { dg-message "note: assigned OpenACC gang loop parallelism" } - do i = 1, N ! { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } + !$acc kernels ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + do i = 1, N ! { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } a(i) = 0 end do !$acc end kernels - !$acc kernels loop ! { dg-message "note: assigned OpenACC gang loop parallelism" } - ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } + !$acc kernels loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } do i = 1, N b(i) = a(N - i + 1) end do !$acc kernels - !$acc loop ! { dg-message "note: assigned OpenACC gang loop parallelism" } - ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } + !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } do i = 1, N b(i) = a(N - i + 1) end do - !$acc loop ! { dg-message "note: assigned OpenACC gang loop parallelism" } - ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } + !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } do i = 1, N c(i) = a(i) * b(i) end do - a(z) = 0 ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } + a(z) = 0 ! { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } - !$acc loop ! { dg-message "note: assigned OpenACC gang loop parallelism" } - ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } + !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } do i = 1, N c(i) = c(i) + a(i) end do - !$acc loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } - ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } + !$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } do i = 1 + 1, N c(i) = c(i) + c(i - 1) end do !$acc end kernels - !$acc kernels ! { dg-bogus "note: assigned OpenACC seq loop parallelism" "TODO" { xfail *-*-* } } - !$acc loop independent ! { dg-message "note: assigned OpenACC gang loop parallelism" } - ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } + !$acc kernels ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } + !$acc loop independent ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } do i = 1, N - !$acc loop independent ! { dg-message "note: assigned OpenACC worker loop parallelism" } + !$acc loop independent ! { dg-message "optimized: assigned OpenACC worker loop parallelism" } do j = 1, N - !$acc loop independent ! { dg-message "note: assigned OpenACC seq loop parallelism" "TODO" { xfail *-*-* } } - ! { dg-warning "insufficient partitioning available to parallelize loop" "TODO" { xfail *-*-* } .-1 } - ! { dg-bogus "note: assigned OpenACC vector loop parallelism" "TODO" { xfail *-*-* } .-2 } + !$acc loop independent ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } + ! { dg-bogus "optimized: assigned OpenACC vector loop parallelism" } do k = 1, N a(1 + mod(i + j + k, N)) & = b(j) & - + f_v (c(k)) ! { dg-message "note: assigned OpenACC vector loop parallelism" "TODO" { xfail *-*-* } .-1 } + + f_v (c(k)) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } end do end do end do !TODO Should the following turn into "gang-single" instead of "parloops"? !TODO The problem is that the first STMT is "if (y <= 4) goto ; else goto ;", thus "parloops". - if (y < 5) then ! { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } - !$acc loop independent ! { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" } + if (y < 5) then ! { dg-message "optimized: beginning .parloops. region in OpenACC .kernels. construct" } + !$acc loop independent ! { dg-message "optimized: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" } do j = 1, N b(j) = f_w (c(j)) end do @@ -101,32 +100,30 @@ program main !$acc end kernels !$acc kernels - !TODO This refers to the "gang-single" "f_g" call. - ! { dg-warning "region contains gang partitoned code but is not gang partitioned" "TODO" { xfail *-*-* } .-2 } - ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" "" { target *-*-* } .+1 } - y = f_g (a(5)) ! { dg-message "note: assigned OpenACC gang worker vector loop parallelism" "TODO" { xfail *-*-* } } - - !$acc loop independent ! { dg-message "note: assigned OpenACC gang loop parallelism" "TODO" { xfail *-*-* } } - ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } - ! { dg-bogus "note: assigned OpenACC gang vector loop parallelism" "TODO" { xfail *-*-* } .-2 } + ! { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" "" { target *-*-* } .+1 } + y = f_g (a(5)) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } + + !$acc loop independent ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } + ! { dg-bogus "optimized: assigned OpenACC gang vector loop parallelism" "" { target *-*-* } .-2 } do j = 1, N - b(j) = y + f_w (c(j)) ! { dg-message "note: assigned OpenACC worker vector loop parallelism" "TODO" { xfail *-*-* } } + b(j) = y + f_w (c(j)) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } end do !$acc end kernels !$acc kernels - y = 3 ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } + y = 3 ! { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } - !$acc loop independent ! { dg-message "note: assigned OpenACC gang worker loop parallelism" "TODO" { xfail *-*-* } } - ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } - ! { dg-bogus "note: assigned OpenACC gang vector loop parallelism" "TODO" { xfail *-*-* } .-2 } + !$acc loop independent ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } + ! { dg-bogus "optimized: assigned OpenACC gang vector loop parallelism" "" { target *-*-* } .-2 } do j = 1, N - b(j) = y + f_v (c(j)) ! { dg-message "note: assigned OpenACC vector loop parallelism" "TODO" { xfail *-*-* } } + b(j) = y + f_v (c(j)) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } end do - z = 2 ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } + z = 2 ! { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } !$acc end kernels - !$acc kernels ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } + !$acc kernels ! { dg-message "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } !$acc end kernels end program main diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index b83ca2d..bc9beba 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -1,6 +1,5 @@ ! { dg-do compile } ! { dg-additional-options "-fdump-tree-original" } -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fdump-tree-convert_oacc_kernels" } program test diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95 index 874c62d..f77bb23 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95 @@ -35,7 +35,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop seq gang ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc loop seq gang ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO @@ -56,11 +56,11 @@ program test !$acc loop worker ! { dg-error "inner loop uses same OpenACC parallelism as containing loop" } DO j = 1,10 ENDDO - !$acc loop gang ! { dg-error "" "TODO" { xfail *-*-* } } + !$acc loop gang DO j = 1,10 ENDDO ENDDO - !$acc loop seq worker ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc loop seq worker ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO !$acc loop gang worker @@ -81,14 +81,14 @@ program test !$acc loop vector ! { dg-error "inner loop uses same OpenACC parallelism as containing loop" } DO j = 1,10 ENDDO - !$acc loop worker ! { dg-error "" "TODO" { xfail *-*-* } } + !$acc loop worker DO j = 1,10 ENDDO - !$acc loop gang ! { dg-error "" "TODO" { xfail *-*-* } } + !$acc loop gang DO j = 1,10 ENDDO ENDDO - !$acc loop seq vector ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc loop seq vector ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO !$acc loop gang vector @@ -101,7 +101,7 @@ program test !$acc loop auto DO i = 1,10 ENDDO - !$acc loop seq auto ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc loop seq auto ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO !$acc loop gang auto ! { dg-error "'auto' conflicts with other OpenACC loop specifiers" } @@ -133,7 +133,7 @@ program test !$acc kernels loop gang(static:*) DO i = 1,10 ENDDO - !$acc kernels loop seq gang ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc kernels loop seq gang ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO @@ -146,7 +146,7 @@ program test !$acc kernels loop worker(num:5) DO i = 1,10 ENDDO - !$acc kernels loop seq worker ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc kernels loop seq worker ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO !$acc kernels loop gang worker @@ -162,7 +162,7 @@ program test !$acc kernels loop vector(length:5) DO i = 1,10 ENDDO - !$acc kernels loop seq vector ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc kernels loop seq vector ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO !$acc kernels loop gang vector @@ -175,7 +175,7 @@ program test !$acc kernels loop auto DO i = 1,10 ENDDO - !$acc kernels loop seq auto ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + !$acc kernels loop seq auto ! { dg-error "'seq' overrides other OpenACC loop specifiers" "TODO" { xfail *-*-* } } DO i = 1,10 ENDDO !$acc kernels loop gang auto ! { dg-error "'auto' conflicts with other OpenACC loop specifiers" } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c index 7cfc364..fd339d2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c @@ -41,6 +41,7 @@ static int state = -1; static acc_device_t acc_device_type; static int acc_device_num; static int num_gangs, num_workers, vector_length; +static int async; static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info) @@ -58,7 +59,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e assert (prof_info->device_type == acc_device_type); assert (prof_info->device_number == acc_device_num); assert (prof_info->thread_id == -1); - assert (prof_info->async == acc_async_sync); + assert (prof_info->async == async); assert (prof_info->async_queue == prof_info->async); assert (prof_info->src_file == NULL); assert (prof_info->func_name == NULL); @@ -154,8 +155,10 @@ int main() acc_device_num = acc_get_device_num (acc_device_type); assert (state == 0); - /* Parallelism dimensions: compiler/runtime decides. */ STATE_OP (state, = 0); + /* Implicit async. */ + async = acc_async_noval; + /* Parallelism dimensions: compiler/runtime decides. */ num_gangs = num_workers = vector_length = 0; { #define N 100 @@ -175,8 +178,10 @@ int main() #undef N } - /* Parallelism dimensions: literal. */ STATE_OP (state, = 0); + /* Explicit async: without argument. */ + async = acc_async_noval; + /* Parallelism dimensions: literal. */ num_gangs = 30; num_workers = 3; vector_length = 5; @@ -184,6 +189,7 @@ int main() #define N 100 int x[N]; #pragma acc kernels \ + async \ num_gangs (30) num_workers (3) vector_length (5) /* { dg-prune-output "using vector_length \\(32\\), ignoring 5" } */ { @@ -200,8 +206,10 @@ int main() #undef N } - /* Parallelism dimensions: variable. */ STATE_OP (state, = 0); + /* Explicit async: variable. */ + async = 123; + /* Parallelism dimensions: variable. */ num_gangs = 22; num_workers = 5; vector_length = 7; @@ -209,6 +217,7 @@ int main() #define N 100 int x[N]; #pragma acc kernels \ + async (async) \ num_gangs (num_gangs) num_workers (num_workers) vector_length (vector_length) /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */ { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index 601e543..45d786d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -1,4 +1,3 @@ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ #undef NDEBUG @@ -12,14 +11,14 @@ int main() #pragma acc kernels { - int c = 234; /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + int c = 234; /* { dg-warning "optimized: beginning .gang-single. region in OpenACC .kernels. construct" } */ -#pragma acc loop independent gang /* { dg-warning "note: assigned OpenACC gang loop parallelism" } */ - /* { dg-warning "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 17 } */ +#pragma acc loop independent gang /* { dg-warning "optimized: assigned OpenACC gang loop parallelism" } */ + /* { dg-warning "optimized: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 16 } */ for (int i = 0; i < N; ++i)