From patchwork Thu Mar 17 08:04:53 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 52027 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 6D2123952016 for ; Thu, 17 Mar 2022 08:05:22 +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 D0A033857C4A for ; Thu, 17 Mar 2022 08:05:00 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org D0A033857C4A Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.90,188,1643702400"; d="scan'208,223";a="73241223" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 17 Mar 2022 00:05:00 -0800 IronPort-SDR: myueAkbe0QBxedXWhQi+t8pNH6PtgPHRG86df/pfXEtLAGO0gthRGh46gudvU2AXe6CyF/xASe 2GwDdcI1WbBpRqQ48PfgHtYCtMHWhYNE8bQ26/4988h8WihHOwjbCOz6VQa7J8LhbmlKTQbjph A9IsAw142zzMcUQ5lWtouTreVnPVcBcro1FKsLueVx/CIMDJK/Fu5JwvGg9jJw3LfDhX6G/o0a P+VRd1/rcFMz+MYKA5zJxTgWPsv2CyZI9aM99GsDC42gWjLu5J8TP2vg59QyItq4i6cxukNH0w Po4= From: Thomas Schwinge To: Subject: Enhance further testcases to verify Openacc 'kernels' decomposition References: <877dqodhi1.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Thu, 17 Mar 2022 09:04:53 +0100 Message-ID: <87pmmlknru.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_MSPIKE_H2, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: Kwok Cheung Yeung Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! Pushed to master branch commit c43cb355f25dd22133d15819bd6ec03d3d3939fd "Enhance further testcases to verify Openacc 'kernels' decomposition", see attached. Kwok, this ought to resolve some testsuite noise from your upcoming og12 branch. Please let me know if there are any more such issues. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From c43cb355f25dd22133d15819bd6ec03d3d3939fd Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 16 Mar 2022 14:19:41 +0100 Subject: [PATCH] Enhance further testcases to verify Openacc 'kernels' decomposition gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-1.c: Enhance. * c-c++-common/goacc/kernels-loop-g.c: Likewise. * c-c++-common/goacc/nesting-1.c: Likewise. * gcc.dg/goacc/nested-function-1.c: Likewise. * gfortran.dg/goacc/common-block-3.f90: Likewise. * gfortran.dg/goacc/nested-function-1.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Enhance. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. --- .../c-c++-common/goacc-gomp/nesting-1.c | 7 ++- .../c-c++-common/goacc/kernels-loop-g.c | 3 ++ gcc/testsuite/c-c++-common/goacc/nesting-1.c | 18 +++++-- .../gcc.dg/goacc/nested-function-1.c | 22 ++++++++ .../gfortran.dg/goacc/common-block-3.f90 | 18 +++++-- .../gfortran.dg/goacc/nested-function-1.f90 | 10 ++++ .../acc_prof-kernels-1.c | 19 +++++-- .../kernels-loop-g.c | 3 ++ .../testsuite/libgomp.oacc-fortran/if-1.f90 | 51 ++++++++++++++++++- 9 files changed, 134 insertions(+), 17 deletions(-) diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c index 39b92712b31..51c5e359f29 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c @@ -1,3 +1,5 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } + /* { dg-additional-options "-fopt-info-omp-note" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } @@ -21,10 +23,11 @@ void f_acc_kernels (void) { #pragma acc kernels - /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 } - { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-2 } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 } */ { int i; + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} "" { target c } .+3 } + { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} "" { target c++ } .+1 } */ #pragma omp atomic write i = 0; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c index 73b469d7061..5bdaa40b02c 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c @@ -1,5 +1,8 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-O2" } */ /* { dg-additional-options "-g" } */ +/*TODO PR100400 { dg-additional-options -fcompare-debug } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-1.c index 83cbff767a4..8c3d1adc785 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -1,3 +1,5 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-all-omp" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } */ @@ -32,11 +34,13 @@ void f_acc_kernels (void) { #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 2; ++i) ; } @@ -63,15 +67,17 @@ f_acc_data (void) } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ ; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 2; ++i) ; } @@ -102,15 +108,17 @@ f_acc_data (void) } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ ; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 2; ++i) ; } diff --git a/gcc/testsuite/gcc.dg/goacc/nested-function-1.c b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c index c34bcb0d601..2e48410b39d 100644 --- a/gcc/testsuite/gcc.dg/goacc/nested-function-1.c +++ b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c @@ -2,6 +2,8 @@ /* See gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 for the Fortran version. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-all-omp" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } @@ -42,6 +44,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(num:local_arg) worker(local_arg) vector(local_arg) \ wait async(local_arg) + /* { dg-note {OpenACC 'kernels' decomposition: variable 'local_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'local_arg' made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'local_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (local_i = 0; local_i < N; ++local_i) @@ -61,6 +68,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(static:local_arg) worker(local_arg) vector(local_arg) \ wait(local_arg, local_arg + 1, local_arg + 2) async + /* { dg-note {OpenACC 'kernels' decomposition: variable 'local_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'local_arg' already made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'local_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (local_i = 0; local_i < N; ++local_i) @@ -87,6 +99,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \ wait async(nonlocal_arg) + /* { dg-note {OpenACC 'kernels' decomposition: variable 'nonlocal_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'nonlocal_arg' made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'nonlocal_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i) @@ -106,6 +123,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \ wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async + /* { dg-note {OpenACC 'kernels' decomposition: variable 'nonlocal_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'nonlocal_arg' already made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'nonlocal_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i) diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 index 9dbfa4cd2f0..6f08d7eb8d5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 @@ -1,5 +1,7 @@ ! { dg-options "-fopenacc -fdump-tree-omplower" } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } @@ -28,7 +30,11 @@ program main a(i) = b(i) + c end do !$acc kernels ! { dg-line l2 } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2 } + ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l2 } + ! { dg-note {variable 'c\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l2 } ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l2 } + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 1, n x(i) = y(i) + c end do @@ -39,10 +45,14 @@ end program main ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } } ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:i \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:c \\\[len: 4\\\]\\)" 1 "omplower" } } ! Expecting no mapping of un-referenced common-blocks variables diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 index 50fd0c82e14..c631f90b27f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 @@ -1,6 +1,8 @@ ! Exercise nested function decomposition, gcc/tree-nested.c. ! See gcc/testsuite/gcc.dg/goacc/nested-function-1.c for the C version. +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-all-omp" } ! { dg-additional-options "--param=openacc-privatization=noisy" } @@ -44,6 +46,8 @@ contains !$acc kernels loop & !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) & !$acc wait async(local_arg) ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } @@ -65,6 +69,8 @@ contains !$acc kernels loop & !$acc gang(static:local_arg) worker(local_arg) vector(local_arg) & !$acc wait(local_arg, local_arg + 1, local_arg + 2) async ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } @@ -95,6 +101,8 @@ contains !$acc kernels loop & !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait async(nonlocal_arg) ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } @@ -116,6 +124,8 @@ contains !$acc kernels loop & !$acc gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } 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 c82a7edbfa0..2c853971474 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 @@ -1,5 +1,7 @@ /* Test dispatch of events to callbacks. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ @@ -74,7 +76,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e assert (prof_info->device_type == acc_device_type); assert (prof_info->device_number == acc_device_num); assert (prof_info->thread_id == -1); - assert (prof_info->async == acc_async_sync); + assert (prof_info->async == acc_async_noval); assert (prof_info->async_queue == prof_info->async); assert (prof_info->src_file == NULL); assert (prof_info->func_name == NULL); @@ -181,10 +183,13 @@ int main() #define N 100 int x[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */ { + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) x[i] = i * i; } @@ -208,11 +213,14 @@ int main() int x[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ \ num_gangs (30) num_workers (3) vector_length (5) - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ /* { dg-warning {using 'vector_length \(32\)', ignoring 5} {} { target { __OPTIMIZE__ && openacc_nvidia_accel_selected } } l_compute$c_compute } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */ { + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) x[i] = i * i; } @@ -236,11 +244,14 @@ int main() int x[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ \ num_gangs (num_gangs) num_workers (num_workers) vector_length (vector_length) - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */ /* { dg-warning {using 'vector_length \(32\)', ignoring runtime setting} {} { target { __OPTIMIZE__ && openacc_nvidia_accel_selected } } l_compute$c_compute } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute } { dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */ { + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) x[i] = i * i; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c index 88258be16bd..e513946ea71 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c @@ -1,3 +1,6 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-g" } */ +/*TODO PR100400 { dg-additional-options -fcompare-debug } */ #include "kernels-loop.c" diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 index 3c4d9a6efb7..c6d67647d4a 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 @@ -1,6 +1,8 @@ ! { dg-do run } ! { dg-additional-options "-cpp" } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-note-omp" } ! { dg-additional-options "-foffload=-fopt-info-note-omp" } @@ -490,6 +492,9 @@ program main a(:) = 4.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -513,6 +518,9 @@ program main a(:) = 16.0 !$acc kernels if (0 == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -530,6 +538,9 @@ program main a(:) = 8.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (one == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -553,6 +564,9 @@ program main a(:) = 22.0 !$acc kernels if (zero == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -570,6 +584,9 @@ program main a(:) = 16.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (.TRUE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -593,6 +610,9 @@ program main a(:) = 76.0 !$acc kernels if (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -612,6 +632,9 @@ program main nn = 1 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (nn == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -637,6 +660,9 @@ program main nn = 0 !$acc kernels if (nn == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -656,6 +682,9 @@ program main nn = 1 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -681,6 +710,9 @@ program main nn = 0; !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -698,6 +730,9 @@ program main a(:) = 91.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -715,6 +750,9 @@ program main a(:) = 43.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (one == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -738,6 +776,9 @@ program main a(:) = 87.0 !$acc kernels if (one == 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -818,7 +859,10 @@ program main !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - !$acc kernels present (a(1:N)) + !$acc kernels present (a(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N b(i) = a(i) end do @@ -862,7 +906,10 @@ program main !$acc data copyout (b(1:N)) if (1 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - !$acc kernels present (a(1:N)) present (b(1:N)) + !$acc kernels present (a(1:N)) present (b(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N b(i) = a(i) end do -- 2.34.1