From patchwork Thu Jan 31 23:59:30 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1034527 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-495057-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="w6+DK2oi"; 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 43rHM02J7Mz9s3l for ; Fri, 1 Feb 2019 11:01:12 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:mime-version:content-type; q=dns; s=default; b=Zy96QOdFhzesJ2OKgUXDqTfSxFF4qSKOzy3eeZ2FWexuY89ZXu h1AP8fGEIzf5SkfkKfr3mpx3h8inZYQAPqE85eTgPQ71M5yi+5AJ8PzTxDrMPOoT fRc5dvyTUREbhDFvcsbohwXENlpUbX0c/SzoWKmQzGYz+xdic2BEbkAOc= 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:from :to:cc:subject:date:message-id:mime-version:content-type; s= default; bh=XzcjyDuwIpfN2VIXMTpFFFYoiJ4=; b=w6+DK2oi+RBKgksIi76q teCAYBveN50xE/XMZR9GPkM4080wlqciA+1TI0XwHmSK6LXUIuJoFI1eq7UxsqQa mzlDZZ8U5+DJvOIAaQc8cHlmP+E14DdSm/MP71Q6PRj7U9n0velgiUJqK9+2yUMu 9j+lZpO421PbgJgXPo7+Qno= Received: (qmail 12676 invoked by alias); 1 Feb 2019 00:00:41 -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 11948 invoked by uid 89); 1 Feb 2019 00:00:31 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, TIME_LIMIT_EXCEEDED autolearn=unavailable version=3.3.2 spammy=visits, representative, transferring, rep 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; Fri, 01 Feb 2019 00:00:18 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-03.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gpMFh-0001ub-V1 from Thomas_Schwinge@mentor.com ; Thu, 31 Jan 2019 16:00:15 -0800 Received: from hertz.schwinge.homeip.net (137.202.0.90) by SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Fri, 1 Feb 2019 00:00:07 +0000 From: Thomas Schwinge To: CC: =?utf-8?b?R2VyZ8O2?= Barany , Tom de Vries Subject: [og8] OpenACC 'kernels' construct changes: splitting of the construct into several regions User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/25.2.2 (x86_64-pc-linux-gnu) Date: Fri, 1 Feb 2019 00:59:30 +0100 Message-ID: MIME-Version: 1.0 Hi! I've just pushed the attached nine patches to openacc-gcc-8-branch: OpenACC 'kernels' construct changes: splitting of the construct into several regions. There's more work to be done there, and we're aware of a number of TODO items, but nevertheless: it's a good first step. (Tom, CCed you just for your information, as you've been working on the OpenACC 'kernels' construct before.) Grüße Thomas From 53f61640a510d13537709239d523c283881f0755 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 23 Jan 2019 02:40:08 -0800 Subject: [PATCH 9/9] Make new OpenACC kernels conversion the default; adjust and add tests 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/dtype-1.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/routine-1.c: Likewise. * c-c++-common/goacc/uninit-dim-clause.c: Likewise. * gfortran.dg/goacc/dtype-1.f95: Likewise. * gfortran.dg/goacc/kernels-conversion.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise. * testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise. * testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise. * testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90: Likewise. --- gcc/ChangeLog.openacc | 4 + gcc/c-family/ChangeLog.openacc | 4 + gcc/c-family/c.opt | 2 +- gcc/doc/invoke.texi | 2 +- gcc/fortran/ChangeLog.openacc | 4 + gcc/fortran/lang.opt | 2 +- gcc/testsuite/ChangeLog.openacc | 40 ++++ .../goacc/classify-kernels-unparallelized.c | 7 +- .../c-c++-common/goacc/classify-kernels.c | 2 +- .../c-c++-common/goacc/classify-parallel.c | 2 +- .../c-c++-common/goacc/classify-routine.c | 2 +- gcc/testsuite/c-c++-common/goacc/dtype-1.c | 8 +- .../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 | 1 - .../c-c++-common/goacc/loop-2-kernels.c | 14 +- ...kernels-conditional-loop-independent_seq.c | 129 +++++++++++ .../note-parallelism-1-kernels-loop-auto.c | 126 +++++++++++ ...rallelism-1-kernels-loop-independent_seq.c | 126 +++++++++++ .../goacc/note-parallelism-1-kernels-loops.c | 47 ++++ ...note-parallelism-1-kernels-straight-line.c | 82 +++++++ ...e-parallelism-combined-kernels-loop-auto.c | 121 +++++++++++ ...sm-combined-kernels-loop-independent_seq.c | 121 +++++++++++ ...kernels-conditional-loop-independent_seq.c | 204 ++++++++++++++++++ .../note-parallelism-kernels-loop-auto.c | 138 ++++++++++++ ...parallelism-kernels-loop-independent_seq.c | 138 ++++++++++++ .../goacc/note-parallelism-kernels-loops.c | 50 +++++ .../c-c++-common/goacc/note-parallelism.c | 3 +- gcc/testsuite/c-c++-common/goacc/routine-1.c | 2 +- .../c-c++-common/goacc/uninit-dim-clause.c | 6 +- gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 | 8 +- .../gfortran.dg/goacc/kernels-conversion.f95 | 7 +- .../gfortran.dg/goacc/kernels-decompose-1.f95 | 1 - .../gfortran.dg/goacc/kernels-tree.f95 | 1 - libgomp/ChangeLog.openacc | 16 ++ .../acc_prof-kernels-1.c | 17 +- .../avoid-offloading-1.c | 18 +- .../avoid-offloading-2.c | 17 +- .../avoid-offloading-3.c | 14 +- .../kernels-decompose-1.c | 3 +- .../libgomp.oacc-fortran/avoid-offloading-1.f | 18 +- .../libgomp.oacc-fortran/avoid-offloading-2.f | 18 +- .../libgomp.oacc-fortran/avoid-offloading-3.f | 15 +- .../initialize_kernels_loops.f90 | 1 - 44 files changed, 1494 insertions(+), 58 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 diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc index 3ef97adef47..433653b2b38 100644 --- a/gcc/ChangeLog.openacc +++ b/gcc/ChangeLog.openacc @@ -1,3 +1,7 @@ +2019-01-31 Thomas Schwinge + + * doc/invoke.texi (-fopenacc-kernels): Update. + 2019-01-31 Thomas Schwinge Gergö Barany diff --git a/gcc/c-family/ChangeLog.openacc b/gcc/c-family/ChangeLog.openacc index 5b60c3a0dee..39d51c53808 100644 --- a/gcc/c-family/ChangeLog.openacc +++ b/gcc/c-family/ChangeLog.openacc @@ -1,3 +1,7 @@ +2019-01-31 Thomas Schwinge + + * c.opt (fopenacc-kernels): Default to "split". + 2019-01-31 Thomas Schwinge * c.opt (fopenacc-kernels): New flag. diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt index 12f8f55c50f..7b7a90e938a 100644 --- a/gcc/c-family/c.opt +++ b/gcc/c-family/c.opt @@ -1618,7 +1618,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 3bbeb8c6839..569c02d1e96 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -2165,9 +2165,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/ChangeLog.openacc b/gcc/fortran/ChangeLog.openacc index acb2177f22f..2715d7aa195 100644 --- a/gcc/fortran/ChangeLog.openacc +++ b/gcc/fortran/ChangeLog.openacc @@ -1,3 +1,7 @@ +2019-01-31 Thomas Schwinge + + * lang.opt (fopenacc-kernels): Default to "split". + 2019-01-31 Thomas Schwinge * lang.opt (fopenacc-kernels): New flag. diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt index b3c9cdb425f..19199aeccbc 100644 --- a/gcc/fortran/lang.opt +++ b/gcc/fortran/lang.opt @@ -643,7 +643,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/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc index 3b4a9c8370d..5dd6d7d656c 100644 --- a/gcc/testsuite/ChangeLog.openacc +++ b/gcc/testsuite/ChangeLog.openacc @@ -1,3 +1,43 @@ +2019-01-31 Thomas Schwinge + Gergö Barany + + * 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/dtype-1.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/routine-1.c: Likewise. + * c-c++-common/goacc/uninit-dim-clause.c: Likewise. + * gfortran.dg/goacc/dtype-1.f95: Likewise. + * gfortran.dg/goacc/kernels-conversion.f95: Likewise. + * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. + * gfortran.dg/goacc/kernels-tree.f95: Likewise. + 2019-01-31 Thomas Schwinge Gergö Barany 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 64467774037..61caa2df236 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 "note: assigned OpenACC seq loop parallelism" } */ for (unsigned int i = 0; i < N; i++) + /* 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 c59a65e1d0f..eb78c8f5e40 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" } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index b345c225aea..baa0bce004d 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 5ca2ec9c603..094c9a760fa 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/dtype-1.c b/gcc/testsuite/c-c++-common/goacc/dtype-1.c index 6dd6ebd8ae1..ae3de574dfc 100644 --- a/gcc/testsuite/c-c++-common/goacc/dtype-1.c +++ b/gcc/testsuite/c-c++-common/goacc/dtype-1.c @@ -96,11 +96,13 @@ test () /* { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ wait\\(10\\) vector_length\\(10\\) num_workers\\(10\\) num_gangs\\(10\\) async\\(10\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(3\\) vector_length\\(128\\) num_workers\\(300\\) num_gangs\\(300\\) async\\(3\\) \\\] wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\) num_gangs\\(1\\) async\\(1\\)" 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ wait\\(-1\\) async\\(-1\\) \\\]" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\) num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ wait\\(-1\\) async\\(-1\\) \\\]" 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ wait\\(1\\) async\\(1\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ wait\\(1\\) async\\(1\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(2\\) async\\(2\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(2\\) async\\(2\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */ + +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ wait\\(1\\) async\\(1\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */ /* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\.0\\) private\\(i1\\)" 1 "omplower" } } */ 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 e17b5dd1107..9920b4fd175 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 ea7eec997fb..8cb63f00444 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 b5d58c37200..7255d692e34 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". */ 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 2608c128b43..8b5c5e75c35 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c @@ -35,7 +35,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++) { } @@ -61,7 +61,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 @@ -90,7 +90,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 @@ -103,7 +103,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 @@ -145,7 +145,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 @@ -161,7 +161,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 @@ -174,7 +174,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 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 00000000000..a81d3559daf --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 00000000000..22ac5399a9d --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto worker + for (y = 0; y < 10; y++) +#pragma acc loop auto vector + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 00000000000..a436cd3f007 --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#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 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#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 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#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 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: 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 00000000000..e8b994b5be0 --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + { + for (x = 0; x < 10; x++) /* { dg-message "note: 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 00000000000..8e40f6217be --- /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 /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */ + { + x = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + y = x < 10; + z = x++; + ; + + y = 0; + z = y < 10; + x -= f_g (y++); /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + ; + + x = f_w (0); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + z = f_v (x < 10); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + y -= f_s (x++); /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + x++; + y = 0; + y += f_v (y < 10); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + y++; + z = 0; + y += f_w (z < 10); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + z++; + ; + + x = 0; + y *= f_g ( /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + f_w (x < 10) /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + + f_g (x < 10) /* { dg-message "note: 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 00000000000..0254036d7af --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto worker + for (y = 0; y < 10; y++) +#pragma acc loop auto vector + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 00000000000..83602a9414d --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: 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 00000000000..e12e0fdae52 --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent worker + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang worker + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent worker vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang worker vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { 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" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { 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" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { 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" "" { 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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { 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" "" { 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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { 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" "" { 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 "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: 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 00000000000..d52b2e860c2 --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto worker + for (y = 0; y < 10; y++) +#pragma acc loop auto vector + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 00000000000..661f7122a2c --- /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 "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "note: 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 "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#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 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#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 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#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 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: 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 00000000000..7587d9d2962 --- /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 "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: 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 7548fb72d14..c514fae2574 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/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index 738957501b9..7cf5506e41a 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -91,7 +91,7 @@ extern void nohost (void); int main () { -#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32) +#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32) /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */ { gang (); worker (); 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 72aacd70f79..6bc35d77765 100644 --- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c +++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c @@ -21,12 +21,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/dtype-1.f95 b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 index 460922a35df..b716d450abf 100644 --- a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 @@ -175,13 +175,13 @@ end subroutine sr5b ! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ async\\(10\\) wait\\(10\\) num_gangs\\(10\\) num_workers\\(10\\) vector_length\\(10\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(3\\) wait\\(3\\) num_gangs\\(300\\) num_workers\\(300\\) vector_length\\(128\\) \\\] async\\(1\\) wait\\(1\\) num_gangs\\(1\\) num_workers\\(1\\) vector_length\\(1\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(-1\\) wait\\(-1\\) \\\]" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\) num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ async\\(-1\\) wait\\(-1\\) \\\]" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) \\\]" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) \\\]" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia\\) \\\[ async\\(2\\) wait\\(2\\) \\\]" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia\\) \\\[ async\\(2\\) wait\\(2\\) \\\]" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(1\\) wait\\(1\\) \\\] async\\(-1\\) wait\\(-1\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(1\\) wait\\(1\\) \\\] async\\(-1\\) wait\\(-1\\)" 1 "omplower" } } ! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\) private\\(i1\\.1\\)" 1 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 index 6604727cf13..4672d15572e 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 520bf034ac6..8173c3651e1 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". diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index b83ca2d8f06..bc9bebac969 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/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc index 27ce3434e49..96908d1a305 100644 --- a/libgomp/ChangeLog.openacc +++ b/libgomp/ChangeLog.openacc @@ -1,5 +1,21 @@ 2019-01-31 Thomas Schwinge + * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: + Update. + * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: + Likewise. + * testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise. + * testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise. + * testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise. + * testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: New file. 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 24b5718ced6..1a5b5fbaa6d 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 @@ -34,6 +34,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; void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info) { @@ -50,7 +51,7 @@ void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *event_in 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); @@ -142,8 +143,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 @@ -172,8 +175,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; @@ -181,6 +186,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" } */ { @@ -206,8 +212,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; @@ -215,6 +223,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/avoid-offloading-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c index 72b9ce0ce02..fb57faef054 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c @@ -1,5 +1,7 @@ /* Test that the compiler decides to "avoid offloading". */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + #include int main(void) @@ -7,8 +9,20 @@ int main(void) int x, y; #pragma acc data copyout(x, y) -#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } */ - *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); +#pragma acc kernels + *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + + /* The following will trigger "avoid offloading". */ +#pragma acc kernels + { +#pragma acc loop auto /* { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } */ + /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 18 } */ + /* { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 18 } */ + for (int i = 0; i < x; ++i) + if (x == 0) + x = 1; + ; + } if (x != 33) __builtin_abort(); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c index 9e05d84d792..00912a3d006 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c @@ -1,6 +1,8 @@ /* Test that a user can override the compiler's "avoid offloading" decision. */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + #include int main(void) @@ -19,8 +21,19 @@ int main(void) int x, y; #pragma acc data copyout(x, y) -#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } */ - *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); +#pragma acc kernels + *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + + /* The following will trigger "avoid offloading". */ +#pragma acc kernels + { +#pragma acc loop auto /* { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } */ + /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 30 } */ + /* { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 30 } */ + for (int i = 0; i < x; ++i) + if (x == 0) + x = 1; + } if (x != 33) __builtin_abort(); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c index f186482026e..1bb6b103b4e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c @@ -1,6 +1,8 @@ /* Test that a user can override the compiler's "avoid offloading" decision. */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + /* Override the compiler's "avoid offloading" decision. { dg-additional-options "-foffload-force" } */ @@ -12,7 +14,17 @@ int main(void) #pragma acc data copyout(x, y) #pragma acc kernels - *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); + *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + + /* The following would trigger "avoid offloading". */ +#pragma acc kernels + { +#pragma acc loop auto /* { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } */ + /* { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 22 } */ + for (int i = 0; i < x; ++i) + if (x == 0) + x = 1; + } if (x != 33) __builtin_abort(); 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 601e543495e..bf58a139d3a 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 @@ -15,7 +14,7 @@ int main() int c = 234; /* { dg-warning "note: 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 } */ + /* { dg-warning "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 16 } */ for (int i = 0; i < N; ++i) b[i] = c; diff --git a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f index fb14be19d8d..275c22f90a6 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-1.f @@ -2,6 +2,7 @@ ! { dg-do run } ! { dg-additional-options "-cpp" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! As __OPTIMIZE__ is defined for -O1 and higher, we don't have an (easy) way to ! distinguish -O1 (where we will offload) from -O2 (where we won't offload), so ! for -O1 testing, we expect to abort. @@ -10,16 +11,29 @@ IMPLICIT NONE INCLUDE "openacc_lib.h" + INTEGER :: I INTEGER, VOLATILE :: X LOGICAL :: Y !$ACC DATA COPYOUT(X, Y) -!$ACC KERNELS ! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } - X = 33 +!$ACC KERNELS + X = 33 ! { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } Y = ACC_ON_DEVICE (ACC_DEVICE_HOST); !$ACC END KERNELS !$ACC END DATA + ! The following will trigger "avoid offloading". +!$ACC KERNELS +!$ACC LOOP AUTO ! { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } +! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 27 } +! { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 27 } + DO I = 1, X + IF (X .EQ. 0) THEN + X = 1 + END IF + END DO +!$ACC END KERNELS + IF (X .NE. 33) CALL ABORT #if defined ACC_DEVICE_TYPE_nvidia # if !defined __OPTIMIZE__ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f index 5a064618e51..a83e1739885 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-2.f @@ -2,11 +2,13 @@ ! { dg-do run } ! { dg-additional-options "-cpp" } +! { dg-additional-options "-fopt-info-optimized-omp" } IMPLICIT NONE INCLUDE "openacc_lib.h" INTEGER :: D + INTEGER :: I INTEGER, VOLATILE :: X LOGICAL :: Y @@ -21,12 +23,24 @@ CALL ACC_INIT (D) !$ACC DATA COPYOUT(X, Y) -!$ACC KERNELS ! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } - X = 33 +!$ACC KERNELS + X = 33 ! { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } Y = ACC_ON_DEVICE (ACC_DEVICE_HOST) !$ACC END KERNELS !$ACC END DATA + ! The following will trigger "avoid offloading". +!$ACC KERNELS +!$ACC LOOP AUTO ! { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } +! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 34 } +! { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 34 } + DO I = 1, X + IF (X .EQ. 0) THEN + X = 1 + END IF + END DO +!$ACC END KERNELS + IF (X .NE. 33) CALL ABORT #if defined ACC_DEVICE_TYPE_nvidia IF (Y) CALL ABORT diff --git a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-3.f b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-3.f index 1c09f83b099..f101186ca19 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-3.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/avoid-offloading-3.f @@ -2,6 +2,7 @@ ! { dg-do run } ! { dg-additional-options "-cpp" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! Override the compiler's "avoid offloading" decision. ! { dg-additional-options "-foffload-force" } @@ -9,16 +10,28 @@ INCLUDE "openacc_lib.h" INTEGER :: D + INTEGER :: I INTEGER, VOLATILE :: X LOGICAL :: Y !$ACC DATA COPYOUT(X, Y) !$ACC KERNELS - X = 33 + X = 33 ! { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } Y = ACC_ON_DEVICE (ACC_DEVICE_HOST) !$ACC END KERNELS !$ACC END DATA + ! The following would trigger "avoid offloading". +!$ACC KERNELS +!$ACC LOOP AUTO ! { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } +! { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 26 } + DO I = 1, X + IF (X .EQ. 0) THEN + X = 1 + END IF + END DO +!$ACC END KERNELS + IF (X .NE. 33) CALL ABORT #if defined ACC_DEVICE_TYPE_nvidia IF (Y) CALL ABORT diff --git a/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90 b/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90 index 35e909f8278..e8b4f3abc05 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90 @@ -1,5 +1,4 @@ ! { dg-do run } -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fopt-info-optimized-omp" } subroutine kernel(lo, hi, a, b, c) -- 2.17.1