From patchwork Fri Feb 26 12:21:54 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1444915 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4Dn8165dLbz9sBJ for ; Fri, 26 Feb 2021 23:22:17 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id A94D93842428; Fri, 26 Feb 2021 12:22:14 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id EF0F8397282B for ; Fri, 26 Feb 2021 12:22:07 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org EF0F8397282B Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: 8tMZhFONcStSKUkNhRkX/yERbOVwPPNtFwdRJ1lWTNMsiHN7OSjibixHkMAUYNseDObsNFyjJU c8qVIOXxFl5asINJ6xwdQpqvN7VjPWLTrsKEmU6+6a4tsZ3Lv3a+GfdMkRWnYSVhoQibenbGxV zrmmZ5nZiekW1WlkN/tP5lFeKlotMkjASSeGX3tCnR5XoXTkCmJOB41ztfpeaFQInPPsYPQn8P 4sSIeuqRQ6W4i02gfVmmKcMgpb9MHXLAY6l0edq62nWG4/6gOJ3UKJfi8TARD80NrNdNq7G75x tRk= X-IronPort-AV: E=Sophos;i="5.81,208,1610438400"; d="scan'208";a="58667836" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 26 Feb 2021 04:22:06 -0800 IronPort-SDR: xq4fqjwKJNtbM2Dv4Jpg98JiN7S7K3jYzL24mWLdaJfFqFkFO2MGpxwh0eHIo5xUz4sRspMDMM 8UY4aAo/yCQijN012O8Qe2/XVBbmSH1CHZlScRiJqjz043B97beA/JZ4n66TOQPe3BMbbKQShm Aa7sXKzL+Ky8hAixP3+aRg3TcYsv/+wdDHoZ7ka3eDn+co+bhN1WLTYGVHXmTTrAUgTP6AKymW twGvsab/rVM8PQx4eXpkMMl5MjQLBfw+S6GXZsu4hzzuPY0PvNvOuM5nV1hEQnDKLCoDlsfb0y 0pM= From: Julian Brown To: Subject: [PATCH] openacc: Warnings for strange partitioning choices for parallel regions Date: Fri, 26 Feb 2021 04:21:54 -0800 Message-ID: <20210226122154.5209-1-julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch adds warnings for strange partitioning choices -- specifying either more or fewer partitioning levels on a parallel directive than the enclosed offload region actually uses. We've used a version of this patch on the og8/og9/og10 branches for quite a while. Versions have been posted for mainline submission as part of a larger patch several times, e.g. here: https://gcc.gnu.org/pipermail/gcc-patches/2018-October/507938.html One motivation for committing this patch is it removes an ongoing divergence in a large number of test cases between the og* branches and mainline, namely whether the added warnings are expected as a result of compiling those test cases, or not. The warnings themselves are perhaps slightly aggressive, but are intended to help the user write more efficient code. Tested with offloading to AMD GCN and (separately) to NVPTX. OK (now or for stage 1)? Julian 2021-02-22 Nathan Sidwell Tom de Vries Thomas Schwinge Julian Brown Kwok Cheung Yeung gcc/ (oacc_validate_dims): Emit warnings about strange partitioning choices. gcc/testsuite/ * c-c++-common/goacc/acc-icf.c: Update. * c-c++-common/classify-serial.c: Likewise. * c-c++-common/kernels-decompose-2.c: Likewise. * c-c++-common/goacc/parallel-dims-1.c: Likewise. * c-c++-common/goacc/parallel-reduction.c: Likewise. * c-c++-common/goacc/pr70688.c: Likewise. * c-c++-common/goacc/routine-1.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise. * c-c++-common/goacc/uninit-dim-clause.c: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/routine-4.f90: Likewise. * gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise. * gfortran.dg/goacc/routine-module-mod-1.f90: Likewise. * gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise. * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Add -w. * testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New. * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise. * testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise. * testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise. * testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise. * testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise. * testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise. --- gcc/omp-offload.c | 28 ++++++++++++++ gcc/testsuite/c-c++-common/goacc/acc-icf.c | 4 +- .../c-c++-common/goacc/classify-serial.c | 2 + .../c-c++-common/goacc/kernels-decompose-2.c | 1 + .../c-c++-common/goacc/parallel-dims-1.c | 6 ++- .../c-c++-common/goacc/parallel-reduction.c | 2 +- gcc/testsuite/c-c++-common/goacc/pr70688.c | 6 +-- gcc/testsuite/c-c++-common/goacc/routine-1.c | 6 +-- .../goacc/routine-level-of-parallelism-2.c | 3 ++ .../c-c++-common/goacc/uninit-dim-clause.c | 3 ++ .../gfortran.dg/goacc/classify-serial.f95 | 2 + .../gfortran.dg/goacc/kernels-decompose-2.f95 | 1 + .../gfortran.dg/goacc/parallel-tree.f95 | 3 ++ gcc/testsuite/gfortran.dg/goacc/routine-4.f90 | 4 ++ .../goacc/routine-level-of-parallelism-1.f90 | 4 +- .../goacc/routine-module-mod-1.f90 | 1 + .../goacc/routine-multiple-directives-1.f90 | 4 +- .../gfortran.dg/goacc/uninit-dim-clause.f95 | 3 ++ .../firstprivate-1.c | 2 + .../libgomp.oacc-c-c++-common/loop-auto-1.c | 2 +- .../libgomp.oacc-c-c++-common/loop-g-1.c | 1 + .../libgomp.oacc-c-c++-common/loop-g-2.c | 1 + .../libgomp.oacc-c-c++-common/loop-red-g-1.c | 1 + .../libgomp.oacc-c-c++-common/loop-red-w-1.c | 1 + .../libgomp.oacc-c-c++-common/loop-red-w-2.c | 1 + .../libgomp.oacc-c-c++-common/loop-w-1.c | 1 + .../libgomp.oacc-c-c++-common/loop-warn-1.c | 37 +++++++++++++++++++ .../mode-transitions.c | 21 +++++++++++ .../libgomp.oacc-c-c++-common/parallel-dims.c | 12 ++++-- .../libgomp.oacc-c-c++-common/pr85381-3.c | 2 +- .../private-variables.c | 13 +++++++ .../libgomp.oacc-c-c++-common/reduction-7.c | 8 ++++ .../libgomp.oacc-c-c++-common/routine-g-1.c | 2 + .../libgomp.oacc-c-c++-common/routine-w-1.c | 1 + .../libgomp.oacc-c-c++-common/routine-wv-2.c | 3 +- .../libgomp.oacc-fortran/optional-private.f90 | 3 ++ .../libgomp.oacc-fortran/par-reduction-2-1.f | 4 +- .../libgomp.oacc-fortran/par-reduction-2-2.f | 4 +- .../libgomp.oacc-fortran/parallel-dims.f90 | 3 ++ .../libgomp.oacc-fortran/pr84028.f90 | 2 +- .../private-variables.f90 | 8 ++++ .../libgomp.oacc-fortran/routine-7.f90 | 2 +- 42 files changed, 191 insertions(+), 27 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index ba0937fba94..57be342da97 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -928,6 +928,34 @@ oacc_validate_dims (tree fn, tree attrs, int *dims, int level, unsigned used) pos = TREE_CHAIN (pos); } + bool check = true; +#ifdef ACCEL_COMPILER + check = false; +#endif + if (check + && !lookup_attribute ("oacc kernels", DECL_ATTRIBUTES (fn))) + { + static char const *const axes[] = + /* Must be kept in sync with GOMP_DIM enumeration. */ + { "gang", "worker", "vector" }; + for (ix = level >= 0 ? level : 0; ix != GOMP_DIM_MAX; ix++) + if (dims[ix] < 0) + ; /* Defaulting axis. */ + else if ((used & GOMP_DIM_MASK (ix)) && dims[ix] == 1) + /* There is partitioned execution, but the user requested a + dimension size of 1. They're probably confused. */ + warning_at (DECL_SOURCE_LOCATION (fn), 0, + "region contains %s partitioned code but" + " is not %s partitioned", axes[ix], axes[ix]); + else if (!(used & GOMP_DIM_MASK (ix)) && dims[ix] != 1) + /* The dimension is explicitly partitioned to non-unity, but + no use is made within the region. */ + warning_at (DECL_SOURCE_LOCATION (fn), 0, + "region is %s partitioned but" + " does not contain %s partitioned code", + axes[ix], axes[ix]); + } + bool changed = targetm.goacc.validate_dims (fn, dims, level, used); /* Default anything left to 1 or a partitioned default. */ diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c index 98b536c34d0..8601ace5db1 100644 --- a/gcc/testsuite/c-c++-common/goacc/acc-icf.c +++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c @@ -4,7 +4,7 @@ #pragma acc routine gang int -routine1 (int n) +routine1 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */ { int i; @@ -17,7 +17,7 @@ routine1 (int n) #pragma acc routine gang int -routine2 (int n) +routine2 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */ { int i; diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c index 94ace1b3c20..89b5f5ed9f9 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-serial.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c @@ -15,6 +15,8 @@ extern unsigned int *__restrict c; void SERIAL () { #pragma acc serial loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ +/* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */ +/* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target *-*-* } .-2 } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c index ec0f75c4a61..f9080936f25 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -121,6 +121,7 @@ main () /*TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"? { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part } */ /* { dg-optimized "assigned OpenACC gang worker vector loop parallelism" "" { target *-*-* } l_part$c_part } */ + /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-6 } */ #pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */ /* { dg-message "note: parallelized loop nest in OpenACC 'kernels' region" "" { target *-*-* } l_loop_j$c_loop_j } */ diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c index 57f682f7a0d..6cdbebe0128 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c +++ b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c @@ -3,9 +3,11 @@ void f(int i) { -#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i) +#pragma acc kernels \ + num_gangs(i) num_workers(i) vector_length(i) ; -#pragma acc parallel num_gangs(i) num_workers(i) vector_length(i) +#pragma acc parallel /* { dg-bogus "region is (gang|worker|vector) partitioned" "" { xfail *-*-* } } */ \ + num_gangs(i) num_workers(i) vector_length(i) ; } diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c index d7cc9470127..9a142c4e02d 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c +++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c @@ -6,7 +6,7 @@ main () #pragma acc data copy (dummy) { -#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) /* { dg-warning "gang partitioned" } */ { int v = 5; sum += 10 + v; diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c index 5a236654069..3f5584a503e 100644 --- a/gcc/testsuite/c-c++-common/goacc/pr70688.c +++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c @@ -21,7 +21,7 @@ parallel_reduction () #pragma acc data copy (dummy) { -#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) /* { dg-warning "region is gang partitioned" } */ { int v = 5; sum += 10 + v; @@ -36,11 +36,11 @@ main () { int i, s = 0; -#pragma acc parallel num_gangs (10) copy (s) reduction (+:s) +#pragma acc parallel num_gangs (10) copy (s) reduction (+:s) /* { dg-warning "region is gang partitioned" } */ for (i = 0; i < n; i++) s += i+1; -#pragma acc parallel num_gangs (10) reduction (+:s) copy (s) +#pragma acc parallel num_gangs (10) reduction (+:s) copy (s) /* { dg-warning "region is gang partitioned" } */ for (i = 0; i < n; i++) s += i+1; diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index a75692246b6..b90e2c11c9d 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -1,16 +1,16 @@ #pragma acc routine gang -void gang (void) +void gang (void) /* { dg-warning "partitioned" 3 } */ { } #pragma acc routine worker -void worker (void) +void worker (void) /* { dg-warning "partitioned" 2 } */ { } #pragma acc routine vector -void vector (void) +void vector (void) /* { dg-warning "partitioned" 1 } */ { } diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c index a066f2b9c2b..5f3cc20c7dc 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c @@ -4,6 +4,9 @@ #pragma acc routine gang void g_1 (void) +/* { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } */ +/* { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-2 } */ +/* { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-3 } */ { } #pragma acc routine (g_1) gang 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 827dac71519..9429e5bfa01 100644 --- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c +++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c @@ -5,12 +5,15 @@ void acc_parallel() int i, j, k; #pragma acc parallel num_gangs(i) /* { dg-warning "is used uninitialized" } */ + /* { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } */ ; #pragma acc parallel num_workers(j) /* { dg-warning "is used uninitialized" } */ + /* { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-1 } */ ; #pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized" } */ + /* { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-1 } */ ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 index 51061afd2c6..d7052bacbe8 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 @@ -15,6 +15,8 @@ program main call setup(a, b) !$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } + ! { dg-warning "not gang partitioned" "" { target *-*-* } .-1 } + ! { dg-warning "not vector partitioned" "" { target *-*-* } .-2 } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 index 22f65e5c694..09bc6626030 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 @@ -119,6 +119,7 @@ program main !$acc end kernels !$acc kernels + ! { dg-warning "not gang partitioned" "" { target *-*-* } .-1 } y = f_g (a(5)) ! { dg-line l_part[incr c_part] } !TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"? ! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index e33653bdd78..83aaf699cdb 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -12,6 +12,9 @@ program test !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & !$acc deviceptr(u), private(v), firstprivate(w) + ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } !$acc end parallel end program test diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 index 6714c7b8229..3fb60e71ae0 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 @@ -123,6 +123,7 @@ contains integer, intent (inout) :: a(N) integer :: i + !$acc loop gang worker vector do i = 1, N a(i) = a(i) - a(i) end do @@ -133,6 +134,7 @@ contains integer, intent (inout) :: a(N) integer :: i + !$acc loop worker vector do i = 1, N a(i) = a(i) - a(i) end do @@ -143,6 +145,7 @@ contains integer, intent (inout) :: a(N) integer :: i + !$acc loop vector do i = 1, N a(i) = a(i) - a(i) end do @@ -153,6 +156,7 @@ contains integer, intent (inout) :: a(N) integer :: i + !$acc loop seq do i = 1, N a(i) = a(i) - a(i) end do diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 index 83b8c24b41d..9986c0c8d50 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 @@ -2,8 +2,10 @@ ! with the OpenACC routine directive. The C/C++ counterpart is ! '../../c-c++-common/goacc/routine-level-of-parallelism-2.c'. -subroutine g_1 +subroutine g_1 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" } !$acc routine gang +! { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "worker partitioned" { xfail *-*-* } .-2 } +! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "worker partitioned" { xfail *-*-* } .-3 } end subroutine g_1 subroutine s_1_2a diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 index 23c673fe3bd..ce4466ab581 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 @@ -50,6 +50,7 @@ contains end do end subroutine w_1 + ! { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .+1 } subroutine g_1 implicit none !$acc routine gang diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 index 6e12ee92155..e39f6b58b25 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 @@ -12,14 +12,14 @@ !$ACC ROUTINE(s_2) END SUBROUTINE s_2 - SUBROUTINE v_1 + SUBROUTINE v_1 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" } !$ACC ROUTINE VECTOR !$ACC ROUTINE VECTOR !$ACC ROUTINE(v_1) VECTOR !$ACC ROUTINE VECTOR END SUBROUTINE v_1 - SUBROUTINE v_2 + SUBROUTINE v_2 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" } !$ACC ROUTINE(v_2) VECTOR !$ACC ROUTINE VECTOR !$ACC ROUTINE(v_2) VECTOR diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 index c77d47a3900..b513a56430f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 @@ -5,12 +5,15 @@ subroutine acc_parallel integer :: i, j, k !$acc parallel num_gangs(i) ! { dg-warning "is used uninitialized" } + ! { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } !$acc end parallel !$acc parallel num_workers(j) ! { dg-warning "is used uninitialized" } + ! { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-1 } !$acc end parallel !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized" } + ! { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-1 } !$acc end parallel end subroutine acc_parallel diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c index 689a443ca43..14bc3af4a97 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c @@ -117,6 +117,8 @@ void t4 () arr[i] = 3; #pragma acc parallel firstprivate(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 119 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 119 } */ { #pragma acc loop gang for (i = 0; i < 32; i++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c index 0273c2bddd7..c6ba9ab8acc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c @@ -150,7 +150,7 @@ int gang_1 (int *ary, int size) { clear (ary, size); -#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size) +#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } } */ { #pragma acc loop auto for (int jx = 0; jx < size / 64; jx++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c index 98f02e9840a..5831327d0ae 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c @@ -1,3 +1,4 @@ +/* { dg-additional-options "-w" } */ #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c index 4152a4e6c82..82e8aae88ac 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c @@ -1,3 +1,4 @@ +/* { dg-additional-options "-w" } */ #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c index 7107502e070..2f3a44f3ba7 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c @@ -1,3 +1,4 @@ +/* { dg-additional-options "-w" } */ #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c index 648f89e1668..7344fa8bf00 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c @@ -1,3 +1,4 @@ +/* { dg-additional-options "-w" } */ #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c index f9fcf3703af..d99877ab8a9 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c @@ -1,3 +1,4 @@ +/* { dg-additional-options "-w" } */ #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c index 5fe486f50a1..8731c805b79 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c @@ -16,6 +16,7 @@ int main () #pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \ copyout(workersize) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 17 } */ { #pragma acc loop worker for (unsigned ix = 0; ix < N; ix++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c new file mode 100644 index 00000000000..20a022f2758 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c @@ -0,0 +1,37 @@ + +/* Check warnings about suboptimal partitioning choices. */ + +int main () +{ + int ary[10]; + +#pragma acc parallel copy(ary) num_gangs (1) /* { dg-warning "is not gang partitioned" } */ + { + #pragma acc loop gang + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel copy(ary) num_workers (1) /* { dg-warning "is not worker partitioned" } */ + { + #pragma acc loop worker + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel copy(ary) num_gangs (8) /* { dg-warning "is gang partitioned" } */ + { + #pragma acc loop worker + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel copy(ary) num_workers (8) /* { dg-warning "is worker partitioned" } */ + { + #pragma acc loop gang + for (int i = 0; i < 10; i++) + ary[i] = i; + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c index 4474c127992..f62daf031a0 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c @@ -287,6 +287,7 @@ void t7() int n = 0; #pragma acc parallel copy(n) \ num_gangs(1) num_workers(1) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 288 } */ { n++; } @@ -310,6 +311,7 @@ void t8() #pragma acc parallel copy(arr) \ num_gangs(gangs) num_workers(1) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 312 } */ { int j; #pragma acc loop gang @@ -339,6 +341,7 @@ void t9() #pragma acc parallel copy(arr) \ num_gangs(gangs) num_workers(1) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 342 } */ { int j; #pragma acc loop gang @@ -371,6 +374,7 @@ void t10() #pragma acc parallel copy(arr) \ num_gangs(gangs) num_workers(1) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 375 } */ { int j; #pragma acc loop gang @@ -404,6 +408,7 @@ void t11() #pragma acc parallel copy(arr) \ num_gangs(1024) num_workers(1) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 409 } */ { int j; @@ -442,6 +447,7 @@ void t12() #pragma acc parallel copyout(fizz, buzz, fizzbuzz) \ num_gangs(NUM_GANGS) num_workers(1) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 448 } */ { int j; @@ -488,6 +494,7 @@ void t13() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 495 } */ { int j; #pragma acc loop gang @@ -613,6 +620,7 @@ void t16() #pragma acc parallel copy(n, arr) \ num_gangs(8) num_workers(16) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 621 } */ { int j; #pragma acc loop gang @@ -665,6 +673,7 @@ void t17() #pragma acc parallel copyin(arr_a) copyout(arr_b) \ num_gangs(num_gangs) num_workers(num_workers) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 674 } */ { int j; #pragma acc loop gang @@ -882,6 +891,8 @@ void t21() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 892 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 892 } */ { int j; #pragma acc loop gang @@ -905,6 +916,8 @@ void t22() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 917 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 917 } */ { int j; #pragma acc loop gang @@ -931,6 +944,8 @@ void t23() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 945 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 945 } */ { int j; #pragma acc loop gang @@ -957,6 +972,8 @@ void t24() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 973 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 973 } */ { int j; #pragma acc loop gang @@ -988,6 +1005,7 @@ void t25() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1006 } */ { int j; #pragma acc loop gang @@ -1020,6 +1038,7 @@ void t26() #pragma acc parallel copy(arr) \ num_gangs(8) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1039 } */ { int j; #pragma acc loop gang @@ -1070,6 +1089,8 @@ void t27() #pragma acc parallel copy(n, arr) copyout(ondev) \ num_gangs(ACTUAL_GANGS) num_workers(8) vector_length(32) + /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "gang" { target *-*-* } 1090 } */ + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1090 } */ { int j; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c index cc4c738c1db..003bcac2413 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c @@ -103,7 +103,7 @@ int main () int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; -#pragma acc parallel copy (gangs_actual) \ +#pragma acc parallel copy (gangs_actual) /* { dg-warning "region contains gang partitioned code but is not gang partitioned" } */ \ num_gangs (GANGS) /* { dg-warning "'num_gangs' value must be positive" "" { target c++ } } */ { /* We're actually executing with num_gangs (1). */ @@ -132,7 +132,7 @@ int main () int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; -#pragma acc parallel copy (workers_actual) \ +#pragma acc parallel copy (workers_actual) /* { dg-warning "region contains worker partitioned code but is not worker partitioned" } */ \ num_workers (WORKERS) /* { dg-warning "'num_workers' value must be positive" "" { target c++ } } */ { /* We're actually executing with num_workers (1). */ @@ -161,7 +161,8 @@ int main () int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; -#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \ +#pragma acc parallel copy (vectors_actual) /* { dg-warning "region contains vector partitioned code but is not vector partitioned" } */ \ + /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 164 } */ \ vector_length (VECTORS) /* { dg-warning "'vector_length' value must be positive" "" { target c++ } } */ { /* We're actually executing with vector_length (1), just the GCC nvptx @@ -205,7 +206,7 @@ int main () int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; -#pragma acc parallel copy (gangs_actual) \ +#pragma acc parallel copy (gangs_actual) /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */ \ reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) \ num_gangs (gangs) { @@ -617,6 +618,9 @@ int main () gangs_max = workers_max = vectors_max = INT_MIN; #pragma acc serial copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \ copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) +/* { dg-warning "not gang partitioned" "" { target *-*-* } 619 } */ +/* { dg-warning "not worker partitioned" "" { target *-*-* } 619 } */ +/* { dg-warning "not vector partitioned" "" { target *-*-* } 619 } */ { if (acc_on_device (acc_device_nvidia)) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c index cddbf271906..53fb049f992 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c @@ -1,6 +1,6 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ -/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ +/* { dg-additional-options "-foffload=-fdump-rtl-mach -w" } */ int a; #pragma acc declare create(a) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c index 53f03d17bb2..f0c3447dfef 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c @@ -22,6 +22,8 @@ void local_g_1() arr[i] = 3; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 24 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 24 } */ { int x; @@ -295,6 +297,8 @@ void loop_g_1() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 299 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 299 } */ { #pragma acc loop gang private(x) for (i = 0; i < 32; i++) @@ -320,6 +324,7 @@ void loop_g_2() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 326 } */ { #pragma acc loop gang private(x) for (i = 0; i < 32; i++) @@ -348,6 +353,7 @@ void loop_g_3() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 355 } */ { #pragma acc loop gang private(x) for (i = 0; i < 32; i++) @@ -376,6 +382,7 @@ void loop_g_4() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 384 } */ { #pragma acc loop gang private(x) for (i = 0; i < 32; i++) @@ -408,6 +415,7 @@ void loop_g_5() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 417 } */ { #pragma acc loop gang private(x) for (i = 0; i < 32; i++) @@ -438,6 +446,7 @@ void loop_g_6() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 448 } */ { #pragma acc loop gang private(pt) for (i = 0; i < 32; i++) @@ -559,6 +568,7 @@ void loop_w_1() arr[i] = i; #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 570 } */ { int j; @@ -875,6 +885,8 @@ void parallel_g_1() arr[i] = 3; #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 887 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 887 } */ { #pragma acc loop gang(static:1) for (i = 0; i < 32; i++) @@ -904,6 +916,7 @@ void parallel_g_2() arr[i] = i; #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 918 } */ { #pragma acc loop gang for (i = 0; i < 32; i++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c index c4940b8ad9d..68ae919ea6f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c @@ -14,6 +14,8 @@ void g_np_1() arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 16 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 16 } */ { #pragma acc loop gang reduction(+:res) for (i = 0; i < 1024; i++) @@ -28,6 +30,8 @@ void g_np_1() res = hres = 1; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 32 } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 32 } */ { #pragma acc loop gang reduction(*:res) for (i = 0; i < 12; i++) @@ -52,6 +56,7 @@ void gv_np_1() arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 58 } */ { #pragma acc loop gang vector reduction(+:res) for (i = 0; i < 1024; i++) @@ -76,6 +81,7 @@ void gw_np_1() arr[i] = i; #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 83 } */ { #pragma acc loop gang worker reduction(+:res) for (i = 0; i < 1024; i++) @@ -239,6 +245,7 @@ void v_p_1() #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ private(res) copyout(out) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 246 } */ { #pragma acc loop gang for (j = 0; j < 32; j++) @@ -315,6 +322,7 @@ void w_p_1() #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \ private(res) copyout(out) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 323 } */ { #pragma acc loop gang for (j = 0; j < 32; j++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c index a164f576bc3..8c3b938506f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c @@ -6,6 +6,8 @@ #pragma acc routine gang void __attribute__ ((noinline)) gang (int ary[N]) +/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 8 } */ +/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 8 } */ { #pragma acc loop gang for (unsigned ix = 0; ix < N; ix++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c index acd9884cbd6..bf9228105c5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c @@ -6,6 +6,7 @@ #pragma acc routine worker void __attribute__ ((noinline)) worker (int ary[N]) +/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 8 } */ { #pragma acc loop worker for (unsigned ix = 0; ix < N; ix++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c index 9769ee72430..de167c55c76 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c @@ -42,7 +42,8 @@ int DoWorkVec (int nw) ary[ix][jx] = 0xdeadbeef; printf ("spawning %d ...", nw); fflush (stdout); - + + /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "vector" { target openacc_radeon_accel_selected } 47 } */ #pragma acc parallel num_workers(nw) vector_length (NUM_VECTORS) copy (ary) { WorkVec ((int *)ary, WIDTH, HEIGHT, nw, NUM_VECTORS); diff --git a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 index 0320bbb3bc9..6bc91b7a0bb 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/optional-private.f90 @@ -30,6 +30,8 @@ contains end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 32 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 32 } !$acc loop gang private(x) do i = 1, 32 x = i * 2; @@ -55,6 +57,7 @@ contains end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 59 } !$acc loop gang private(pt) do i = 0, 31 pt%x = i diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f index aa1bb634ba6..ff31116d980 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-1.f @@ -14,7 +14,7 @@ RES2 = 0 !$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32) -!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) +!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" } res1 = res1 + 5 !$ACC ATOMIC @@ -36,7 +36,7 @@ RES2 = 1 !$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32) -!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) +!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" } res1 = res1 * 5 !$ACC ATOMIC diff --git a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f index 5694de1aee8..47c5ff3474d 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f +++ b/libgomp/testsuite/libgomp.oacc-fortran/par-reduction-2-2.f @@ -14,7 +14,7 @@ RES2 = 0 !$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32) -!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) +!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" } res1 = res1 + 5 !$ACC ATOMIC @@ -36,7 +36,7 @@ RES2 = 1 !$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32) -!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) +!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" } res1 = res1 * 5 !$ACC ATOMIC diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 index 1bfcd6ce099..b515c62391b 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 @@ -84,6 +84,9 @@ program main vectors_max = -huge(gangs_max) - 1 ! INT_MIN !$acc serial copy (vectors_actual) & !$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } + ! { dg-warning "not gang partitioned" "" { target *-*-* } 86 } + ! { dg-warning "not worker partitioned" "" { target *-*-* } 86 } + ! { dg-warning "not vector partitioned" "" { target *-*-* } 86 } if (acc_on_device (acc_device_nvidia)) then ! The GCC nvptx back end enforces vector_length (32). ! It's unclear if that's actually permissible here; diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90 index 2b361220bb6..8cb76a93d7f 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr84028.f90 @@ -5,7 +5,7 @@ program foo a = 1 - !$acc parallel num_gangs(1) num_workers(2) + !$acc parallel num_gangs(1) num_workers(2) ! { dg-warning "region is worker partitioned" } if (any(a(1:3,1:3,1:3).ne.1)) STOP 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 index 472a6a14fff..fbff5ccd46e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 @@ -13,6 +13,8 @@ subroutine t1() end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 15 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 15 } !$acc loop gang private(x) do i = 1, 32 x = i * 2; @@ -37,6 +39,7 @@ subroutine t2() end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 41 } !$acc loop gang private(x) do i = 0, 31 x = i * 2; @@ -65,6 +68,7 @@ subroutine t3() end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 70 } !$acc loop gang private(x) do i = 0, 31 x = i * 2; @@ -98,6 +102,7 @@ subroutine t4() end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 104 } !$acc loop gang private(pt) do i = 0, 31 pt%x = i @@ -208,6 +213,7 @@ subroutine t7() end do !$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 215 } !$acc loop gang private(x) do i = 0, 31 !$acc loop worker private(x) @@ -507,6 +513,8 @@ subroutine t14() end do !$acc parallel private(x) copy(arr) num_gangs(n) num_workers(8) vector_length(32) + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 515 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 515 } !$acc loop gang(static:1) do i = 1, n x = i * 2; diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 index 1009f4a81e5..51d52456733 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 @@ -100,7 +100,7 @@ subroutine gang (a) integer, intent (inout) :: a(N) integer :: i - !$acc loop gang + !$acc loop gang worker vector do i = 1, N a(i) = a(i) - i end do