From patchwork Thu Jul 29 07:49:05 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1511111 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=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (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 4Gb2jy1FM9z9sSs for ; Thu, 29 Jul 2021 17:49:41 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id A1E7C393A419 for ; Thu, 29 Jul 2021 07:49:39 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id D10B5385481E for ; Thu, 29 Jul 2021 07:49:18 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org D10B5385481E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 1qClXCE9ClOGnzEnzGHI1V8Nu6u4JcXJa01zU+OxecDtvnHtEaw3mloKflfFBU50zNrUHo+Oa+ tECK7W2XeP1h7c18vCcoGQO553EY+YnKkFE4TZHO4Cdz6HqfjYUC2Jrj9GkVxw3IRyNc2YgPno Rijs3l+8POG5o3GlvhyUWGz1nt8YMUFmocJnd89zLwcF7tLV8RQOjnuf1+Siu7CWLjy7fuNYea g9+GD1gqFm7fE55TJ/XvQpCXnUsnOoUr2jJ8XOrbnJ6Q4I14I7uFuc64ZheQjq1S2WM2mMU8PY oOyFRTr6DOWuNl4dypqB8Dun X-IronPort-AV: E=Sophos;i="5.84,278,1620720000"; d="scan'208,223";a="64046142" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 28 Jul 2021 23:49:17 -0800 IronPort-SDR: 2moupMGxVn9CM/Ab6UNprjzOinFozSpiBuVLEcs2434P9HCQ6WAusMJRY0mQxI+XoqekR0DF7/ 5SlLuShSN40FXh9Vqa2ffd5ORld/DkXlsHYYrtBpZHGwNDZzONO1CIYlu5uL3Vg1LhsO1AidoM FdQAG5aJe0FtpRFhHlmuKgltHePlzwdclNlNwbpcmgO12hsAiR+miU74uyROUk4K9NeFrLE7fd kFACL2RzfiSwOxvnOVhF4VdI5ehiWTSouuoqWoRUUzvQAvyb2r8qT4ZVP+b+9UBqbF6ABsaJ3B PNA= From: Thomas Schwinge To: Julian Brown , Subject: [OpenACC] Extract 'pass_oacc_loop_designation' out of 'pass_oacc_device_lower' (was: [PATCH 1/4] openacc: Middle-end worker-partitioning support) In-Reply-To: <2ef7b2ebaf056858d6484a260c3897f844e2df4a.1614685766.git.julian@codesourcery.com> References: <2ef7b2ebaf056858d6484a260c3897f844e2df4a.1614685766.git.julian@codesourcery.com> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Thu, 29 Jul 2021 09:49:05 +0200 Message-ID: <875ywttv9q.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-06.mgc.mentorg.com (139.181.222.6) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP 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: Jakub Jelinek , Tobias Burnus Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi Julian! On 2021-03-02T04:20:11-0800, Julian Brown wrote: > This patch implements worker-partitioning support in the middle end, > [...] I've first separately pushed the mostly "mechanical changes" re "[OpenACC] Extract 'pass_oacc_loop_designation' out of 'pass_oacc_device_lower'" to master branch in commit 0829ab79d37be6c59072af0c4f54043f7e9d23ea, see attached. A few comments there: > --- a/gcc/omp-offload.c > +++ b/gcc/omp-offload.c > @@ -1367,6 +1368,8 @@ oacc_loop_xform_head_tail (gcall *from, int level) > else if (gimple_call_internal_p (stmt, IFN_GOACC_REDUCTION)) > *gimple_call_arg_ptr (stmt, 3) = replacement; > > + update_stmt (stmt); > + > gsi_next (&gsi); > while (gsi_end_p (gsi)) > gsi = gsi_start_bb (single_succ (gsi_bb (gsi))); > @@ -1391,25 +1394,28 @@ oacc_loop_process (oacc_loop *loop) > [...] > + update_stmt (call); Sneaky. ACK. > /* Main entry point for oacc transformations which run on the device > compiler after LTO, so we know what the target device is at this > point (including the host fallback). */ > > static unsigned int > -execute_oacc_device_lower () > +execute_oacc_loop_designation () This does not just OpenACC loop designation but also includes the general OpenACC offloaded function classification (diagnostics) as well as OpenACC 'nohost' clause handling for OpenACC 'routine', meaning that the "loop designation" name is not totally accurate. But I couldn't easily come up with anything more accurate (or an easy way to split out these things), so I left it at that. (Also, for later, I wonder if not all the 'oacc_loop' stuff could/should move into its own new file 'gcc/omp-oacc-loop.cc'. Also, the tag 'oacc_loop' isn't totally accurate either, for this also deals with OpenACC 'routine' level of parallelism -- maybe 'oacc_lop' instead of 'oacc_loop' etc.) > @@ -2051,10 +2072,36 @@ execute_oacc_device_lower () > free_oacc_loop (l); > } > > + free_oacc_loop (loops); > + > /* Offloaded targets may introduce new basic blocks, which require > dominance information to update SSA. */ > calculate_dominance_info (CDI_DOMINATORS); > > + return 0; > +} I do confirm the manual 'calculate_dominance_info (CDI_DOMINATORS)' necessary in the original state (where this is in the middle of the two "passes"), but given 'TODO_cleanup_cfg' as part of 'todo_flags_finish' for new 'pass_oacc_loop_designation', we no longer need that now, as far as I can tell. So I removed the manual 'calculate_dominance_info (CDI_DOMINATORS)' -- but please do tell if there is a reason to keep it. > namespace { > > +const pass_data pass_data_oacc_loop_designation = > +{ > + GIMPLE_PASS, /* type */ > + "oaccloops", /* name */ > + OPTGROUP_OMP, /* optinfo_flags */ > + TV_NONE, /* tv_id */ > + PROP_cfg, /* properties_required */ > + 0 /* Possibly PROP_gimple_eomp. */, /* properties_provided */ > + 0, /* properties_destroyed */ > + 0, /* todo_flags_start */ > + TODO_update_ssa | TODO_cleanup_cfg > + | TODO_rebuild_alias, /* todo_flags_finish */ > +}; Do you remember why you added 'TODO_rebuild_alias' here? 'pass_oacc_device_lower' doesn't have it, and neither does 'pass_oacc_loop_designation' in your original (2017-11-27) internal gcn/master branch commit 81ee7ef64cdfa47c01f24c79b8ebd03242c9f3eb "Split device-lowering/gimple workers into three passes". So I removed that -- but please do tell if there is a reason to keep it. 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 0829ab79d37be6c59072af0c4f54043f7e9d23ea Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 2 Mar 2021 04:20:11 -0800 Subject: [PATCH] [OpenACC] Extract 'pass_oacc_loop_designation' out of 'pass_oacc_device_lower' This really is a separate step -- and another pass to be added between the two, later on. gcc/ * omp-offload.c (oacc_loop_xform_head_tail, oacc_loop_process): 'update_stmt' after modification. (pass_oacc_loop_designation): New function, extracted out of... (pass_oacc_device_lower): ... this. (pass_data_oacc_loop_designation, pass_oacc_loop_designation) (make_pass_oacc_loop_designation): New * passes.def: Add it. * tree-parloops.c (create_parallel_loop): Adjust. * tree-pass.h (make_pass_oacc_loop_designation): New. gcc/testsuite/ * c-c++-common/goacc/classify-kernels-unparallelized.c: 's%oaccdevlow%oaccloops%g'. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/classify-parallel.c: Likewise. * c-c++-common/goacc/classify-routine-nohost.c: Likewise. * c-c++-common/goacc/classify-routine.c: Likewise. * c-c++-common/goacc/classify-serial.c: Likewise. * c-c++-common/goacc/routine-nohost-1.c: Likewise. * g++.dg/goacc/template.C: Likewise. * gcc.dg/goacc/loop-processing-1.c: Likewise. * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise. * gfortran.dg/goacc/classify-kernels.f95: Likewise. * gfortran.dg/goacc/classify-parallel.f95: Likewise. * gfortran.dg/goacc/classify-routine-nohost.f95: Likewise. * gfortran.dg/goacc/classify-routine.f95: Likewise. * gfortran.dg/goacc/classify-serial.f95: Likewise. * gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: 's%oaccdevlow%oaccloops%g'. * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: Likewise. * testsuite/libgomp.oacc-fortran/routine-nohost-1.f90: Likewise. Co-Authored-By: Julian Brown Co-Authored-By: Kwok Cheung Yeung --- gcc/omp-offload.c | 98 ++++++++++++++----- gcc/passes.def | 1 + .../goacc/classify-kernels-unparallelized.c | 8 +- .../c-c++-common/goacc/classify-kernels.c | 8 +- .../c-c++-common/goacc/classify-parallel.c | 8 +- .../goacc/classify-routine-nohost.c | 22 ++--- .../c-c++-common/goacc/classify-routine.c | 22 ++--- .../c-c++-common/goacc/classify-serial.c | 8 +- .../c-c++-common/goacc/routine-nohost-1.c | 8 +- gcc/testsuite/g++.dg/goacc/template.C | 20 ++-- .../gcc.dg/goacc/loop-processing-1.c | 4 +- .../goacc/classify-kernels-unparallelized.f95 | 8 +- .../gfortran.dg/goacc/classify-kernels.f95 | 8 +- .../gfortran.dg/goacc/classify-parallel.f95 | 8 +- .../goacc/classify-routine-nohost.f95 | 20 ++-- .../gfortran.dg/goacc/classify-routine.f95 | 20 ++-- .../gfortran.dg/goacc/classify-serial.f95 | 8 +- .../goacc/routine-multiple-directives-1.f90 | 34 +++---- gcc/tree-parloops.c | 2 +- gcc/tree-pass.h | 1 + .../libgomp.oacc-c-c++-common/pr85486-2.c | 4 +- .../libgomp.oacc-c-c++-common/pr85486-3.c | 4 +- .../libgomp.oacc-c-c++-common/pr85486.c | 4 +- .../routine-nohost-1.c | 8 +- .../vector-length-128-1.c | 4 +- .../vector-length-128-2.c | 4 +- .../vector-length-128-3.c | 4 +- .../vector-length-128-4.c | 4 +- .../vector-length-128-5.c | 4 +- .../vector-length-128-6.c | 4 +- .../vector-length-128-7.c | 4 +- .../libgomp.oacc-fortran/routine-nohost-1.f90 | 6 +- 32 files changed, 213 insertions(+), 157 deletions(-) diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index bfbb0112e24..d881426ae65 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -1367,6 +1367,7 @@ oacc_loop_xform_head_tail (gcall *from, int level) } else if (gimple_call_internal_p (stmt, IFN_GOACC_REDUCTION)) *gimple_call_arg_ptr (stmt, 3) = replacement; + update_stmt (stmt); gsi_next (&gsi); while (gsi_end_p (gsi)) @@ -1392,25 +1393,28 @@ oacc_loop_process (oacc_loop *loop) gcall *call; for (ix = 0; loop->ifns.iterate (ix, &call); ix++) - switch (gimple_call_internal_fn (call)) - { - case IFN_GOACC_LOOP: + { + switch (gimple_call_internal_fn (call)) { - bool is_e = gimple_call_arg (call, 5) == integer_minus_one_node; - gimple_call_set_arg (call, 5, is_e ? e_mask_arg : mask_arg); - if (!is_e) - gimple_call_set_arg (call, 4, chunk_arg); - } - break; + case IFN_GOACC_LOOP: + { + bool is_e = gimple_call_arg (call, 5) == integer_minus_one_node; + gimple_call_set_arg (call, 5, is_e ? e_mask_arg : mask_arg); + if (!is_e) + gimple_call_set_arg (call, 4, chunk_arg); + } + break; - case IFN_GOACC_TILE: - gimple_call_set_arg (call, 3, mask_arg); - gimple_call_set_arg (call, 4, e_mask_arg); - break; + case IFN_GOACC_TILE: + gimple_call_set_arg (call, 3, mask_arg); + gimple_call_set_arg (call, 4, e_mask_arg); + break; - default: - gcc_unreachable (); - } + default: + gcc_unreachable (); + } + update_stmt (call); + } unsigned dim = GOMP_DIM_GANG; unsigned mask = loop->mask | loop->e_mask; @@ -1912,7 +1916,7 @@ is_sync_builtin_call (gcall *call) point (including the host fallback). */ static unsigned int -execute_oacc_device_lower () +execute_oacc_loop_designation () { tree attrs = oacc_get_fn_attrib (current_function_decl); @@ -1981,6 +1985,8 @@ execute_oacc_device_lower () gcc_unreachable (); } + /* This doesn't belong into 'pass_oacc_loop_designation' conceptually, but + it's a convenient place, so... */ if (is_oacc_routine) { tree attr = lookup_attribute ("omp declare target", @@ -2088,9 +2094,23 @@ execute_oacc_device_lower () free_oacc_loop (l); } - /* Offloaded targets may introduce new basic blocks, which require - dominance information to update SSA. */ - calculate_dominance_info (CDI_DOMINATORS); + free_oacc_loop (loops); + + return 0; +} + +static unsigned int +execute_oacc_device_lower () +{ + tree attrs = oacc_get_fn_attrib (current_function_decl); + + if (!attrs) + /* Not an offloaded function. */ + return 0; + + int dims[GOMP_DIM_MAX]; + for (unsigned i = 0; i < GOMP_DIM_MAX; i++) + dims[i] = oacc_get_fn_dim_size (current_function_decl, i); hash_map adjusted_vars; @@ -2355,8 +2375,6 @@ execute_oacc_device_lower () } } - free_oacc_loop (loops); - return 0; } @@ -2397,6 +2415,36 @@ default_goacc_dim_limit (int ARG_UNUSED (axis)) namespace { +const pass_data pass_data_oacc_loop_designation = +{ + GIMPLE_PASS, /* type */ + "oaccloops", /* name */ + OPTGROUP_OMP, /* optinfo_flags */ + TV_NONE, /* tv_id */ + PROP_cfg, /* properties_required */ + 0 /* Possibly PROP_gimple_eomp. */, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + TODO_update_ssa | TODO_cleanup_cfg, /* todo_flags_finish */ +}; + +class pass_oacc_loop_designation : public gimple_opt_pass +{ +public: + pass_oacc_loop_designation (gcc::context *ctxt) + : gimple_opt_pass (pass_data_oacc_loop_designation, ctxt) + {} + + /* opt_pass methods: */ + virtual bool gate (function *) { return flag_openacc; }; + + virtual unsigned int execute (function *) + { + return execute_oacc_loop_designation (); + } + +}; // class pass_oacc_loop_designation + const pass_data pass_data_oacc_device_lower = { GIMPLE_PASS, /* type */ @@ -2429,6 +2477,12 @@ public: } // anon namespace +gimple_opt_pass * +make_pass_oacc_loop_designation (gcc::context *ctxt) +{ + return new pass_oacc_loop_designation (ctxt); +} + gimple_opt_pass * make_pass_oacc_device_lower (gcc::context *ctxt) { diff --git a/gcc/passes.def b/gcc/passes.def index e2858368b7d..26d86df2f5a 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -183,6 +183,7 @@ along with GCC; see the file COPYING3. If not see INSERT_PASSES_AFTER (all_passes) NEXT_PASS (pass_fixup_cfg); NEXT_PASS (pass_lower_eh_dispatch); + NEXT_PASS (pass_oacc_loop_designation); NEXT_PASS (pass_oacc_device_lower); NEXT_PASS (pass_omp_device_lower); NEXT_PASS (pass_omp_target_link); diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 218f6248062..1d12658790d 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -5,7 +5,7 @@ { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } - { dg-additional-options "-fdump-tree-oaccdevlow" } */ + { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -38,6 +38,6 @@ void KERNELS () /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). - { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 95a150ca9ac..bdf7b4a0641 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -5,7 +5,7 @@ { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } - { dg-additional-options "-fdump-tree-oaccdevlow" } */ + { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -34,6 +34,6 @@ void KERNELS () /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). - { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 230e70c66cd..9056aa69dad 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -4,7 +4,7 @@ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } - { dg-additional-options "-fdump-tree-oaccdevlow" } */ + { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -27,6 +27,6 @@ void PARALLEL () /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). - { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c b/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c index a58482f7f92..99855822011 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c @@ -4,7 +4,7 @@ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } - { dg-additional-options "-fdump-tree-oaccdevlow" } */ + { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -28,14 +28,14 @@ void ROUTINE () { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(nohost worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ /* Check the offloaded function's classification. - { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' has 'nohost' clause" 1 "oaccdevlow" { target c } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' discarded" 1 "oaccdevlow" { target c } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' discarded" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' discarded" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } } + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' has 'nohost' clause" 1 "oaccloops" { target c } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccloops" { target { c++ && offloading_enabled } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' discarded" 1 "oaccloops" { target c } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' discarded" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' discarded" 1 "oaccloops" { target { c++ && offloading_enabled } } } } TODO See PR101551 for 'offloading_enabled' differences. - { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccdevlow" } } - { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccdevlow" } } - { dg-final { scan-tree-dump-not "(?n)void ROUTINE \\(\\)" "oaccdevlow" } } */ + { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccloops" } } + { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccloops" } } + { dg-final { scan-tree-dump-not "(?n)void ROUTINE \\(\\)" "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index cc0ba2b9a7d..f7f0454009b 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -4,7 +4,7 @@ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } - { dg-additional-options "-fdump-tree-oaccdevlow" } */ + { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -29,14 +29,14 @@ void ROUTINE () /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). - { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccdevlow" { target c } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccdevlow" { target c } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } } - { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } } + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccloops" { target c } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops" { target { c++ && offloading_enabled } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccloops" { target c } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccloops" { target { c++ && offloading_enabled } } } } TODO See PR101551 for 'offloading_enabled' differences. - { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c index ae052ae6a1c..f41c141bcd5 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-serial.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c @@ -4,7 +4,7 @@ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } - { dg-additional-options "-fdump-tree-oaccdevlow" } */ + { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -32,6 +32,6 @@ void SERIAL () /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). - { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccloops" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c index c8927416efa..59ebb2bc5a9 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c @@ -1,6 +1,6 @@ /* Test OpenACC 'routine' with 'nohost' clause, valid use. */ -/* { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-fdump-tree-oaccloops" } */ #pragma acc routine nohost int THREE(void) @@ -13,7 +13,7 @@ int THREE(void) #pragma acc routine nohost extern int THREE(void); -/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */ +/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 oaccloops } } */ #pragma acc routine nohost @@ -30,7 +30,7 @@ extern void NOTHING(void); #pragma acc routine (NOTHING) nohost -/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */ +/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 oaccloops } } */ extern float ADD(float, float); @@ -47,4 +47,4 @@ extern float ADD(float, float); #pragma acc routine (ADD) nohost -/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */ +/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 oaccloops } } */ diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C index f34fcfea52d..10d3f446da7 100644 --- a/gcc/testsuite/g++.dg/goacc/template.C +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -1,4 +1,4 @@ -/* { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-fdump-tree-oaccloops" } */ #pragma acc routine nohost template T @@ -156,13 +156,13 @@ main () return b + c; } -/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']+' has 'nohost' clause\.$} 4 oaccdevlow } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = char\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)char' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = int\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)int' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = float\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)float' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = double\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)double' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } } +/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']+' has 'nohost' clause\.$} 4 oaccloops } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = char\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)char' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = int\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)int' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = float\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)float' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = double\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble\(int\)double' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } } TODO See PR101551 for 'offloading_enabled' differences. */ diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c index bd4c07e7d81..78b9aed89be 100644 --- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c +++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -1,5 +1,5 @@ /* Make sure that OpenACC loop processing happens. */ -/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-O2 -fdump-tree-oaccloops" } */ extern int place (); @@ -15,4 +15,4 @@ void vector_1 (int *ary, int size) } } -/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccdevlow" } } */ +/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccloops" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index cb5251a2aeb..3fb48b321f2 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -5,7 +5,7 @@ ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting ! aspects of that functionality. @@ -40,6 +40,6 @@ end program main ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). -! { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index 07aaf065e1d..6c8d298e236 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -5,7 +5,7 @@ ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting ! aspects of that functionality. @@ -36,6 +36,6 @@ end program main ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). -! { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 index a41e0e68b38..ce4c08ff219 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 @@ -4,7 +4,7 @@ ! { dg-additional-options "-O2" } ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting ! aspects of that functionality. @@ -29,6 +29,6 @@ end program main ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). -! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 index 0e06fb9f0ba..07e2063551f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 @@ -4,7 +4,7 @@ ! { dg-additional-options "-O2" } ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting ! aspects of that functionality. @@ -27,13 +27,13 @@ end subroutine ROUTINE ! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 0, 1 0\\), omp declare target \\(nohost worker\\)\\)\\)" 1 "ompexp" } } ! Check the offloaded function's classification. -! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' discarded" 1 "oaccdevlow" { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' discarded" 1 "oaccdevlow" { target offloading_enabled } } } -! { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccdevlow" } } -! { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccdevlow" } } -! { dg-final { scan-tree-dump-not "(?n)void routine \\(\\)" "oaccdevlow" { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-not "(?n)void routine_ \\(\\)" "oaccdevlow" { target offloading_enabled } } } +! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' discarded" 1 "oaccloops" { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' discarded" 1 "oaccloops" { target offloading_enabled } } } +! { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccloops" } } +! { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccloops" } } +! { dg-final { scan-tree-dump-not "(?n)void routine \\(\\)" "oaccloops" { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-not "(?n)void routine_ \\(\\)" "oaccloops" { target offloading_enabled } } } !TODO See PR101551 for 'offloading_enabled' differences. diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index 92d3243cdcf..b065ccadacd 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -4,7 +4,7 @@ ! { dg-additional-options "-O2" } ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting ! aspects of that functionality. @@ -28,13 +28,13 @@ end subroutine ROUTINE ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). -! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' not discarded" 1 "oaccdevlow" { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' not discarded" 1 "oaccdevlow" { target offloading_enabled } } } -! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)void routine \\(\\)" 1 "oaccdevlow" { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-times "(?n)void routine_ \\(\\)" 1 "oaccdevlow" { target offloading_enabled } } } +! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' not discarded" 1 "oaccloops" { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' not discarded" 1 "oaccloops" { target offloading_enabled } } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)void routine \\(\\)" 1 "oaccloops" { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-times "(?n)void routine_ \\(\\)" 1 "oaccloops" { target offloading_enabled } } } !TODO See PR101551 for 'offloading_enabled' differences. diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 index 6dcb1b170f8..f5cb3fe50c5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 @@ -4,7 +4,7 @@ ! { dg-additional-options "-O2" } ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting ! aspects of that functionality. @@ -32,6 +32,6 @@ end program main ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). -! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccloops" } } 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 44ef4533f04..42bcb0e8d63 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 @@ -1,6 +1,6 @@ ! Check for valid cases of multiple OpenACC 'routine' directives. -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } !TODO See PR101551 for 'offloading_enabled' differences. ! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting @@ -11,32 +11,32 @@ !$ACC ROUTINE(s_1) SEQ !$ACC ROUTINE SEQ END SUBROUTINE s_1 - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE s_1_nh !$ACC ROUTINE(s_1_nh) NOHOST !$ACC ROUTINE(s_1_nh) SEQ NOHOST !$ACC ROUTINE NOHOST SEQ END SUBROUTINE s_1_nh - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE s_2 !$ACC ROUTINE !$ACC ROUTINE SEQ !$ACC ROUTINE(s_2) END SUBROUTINE s_2 - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE s_2_nh !$ACC ROUTINE NOHOST !$ACC ROUTINE NOHOST SEQ !$ACC ROUTINE(s_2_nh) NOHOST END SUBROUTINE s_2_nh - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE v_1 !$ACC ROUTINE VECTOR @@ -45,8 +45,8 @@ !$ACC ROUTINE VECTOR ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 } END SUBROUTINE v_1 - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE v_1_nh !$ACC ROUTINE NOHOST VECTOR @@ -55,8 +55,8 @@ !$ACC ROUTINE VECTOR NOHOST ! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 } END SUBROUTINE v_1_nh - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE v_2 !$ACC ROUTINE(v_2) VECTOR @@ -64,8 +64,8 @@ !$ACC ROUTINE(v_2) VECTOR ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } END SUBROUTINE v_2 - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE v_2_nh !$ACC ROUTINE(v_2_nh) VECTOR NOHOST @@ -73,8 +73,8 @@ !$ACC ROUTINE(v_2_nh) NOHOST VECTOR ! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } END SUBROUTINE v_2_nh - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } } - ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } } + ! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } } SUBROUTINE sub_1 IMPLICIT NONE diff --git a/gcc/tree-parloops.c b/gcc/tree-parloops.c index bb547572653..4d40c96df7d 100644 --- a/gcc/tree-parloops.c +++ b/gcc/tree-parloops.c @@ -2867,7 +2867,7 @@ create_parallel_loop (class loop *loop, tree loop_fn, tree data, /* Emit GIMPLE_OMP_FOR. */ if (oacc_kernels_p) /* Parallelized OpenACC kernels constructs use gang parallelism. See also - omp-offload.c:execute_oacc_device_lower. */ + omp-offload.c:execute_oacc_loop_designation. */ t = build_omp_clause (loc, OMP_CLAUSE_GANG); else { diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h index 1f5b1370a95..5484ad5eac7 100644 --- a/gcc/tree-pass.h +++ b/gcc/tree-pass.h @@ -424,6 +424,7 @@ extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt); extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt); extern gimple_opt_pass *make_pass_expand_omp_ssa (gcc::context *ctxt); extern gimple_opt_pass *make_pass_omp_target_link (gcc::context *ctxt); +extern gimple_opt_pass *make_pass_oacc_loop_designation (gcc::context *ctxt); extern gimple_opt_pass *make_pass_oacc_device_lower (gcc::context *ctxt); extern gimple_opt_pass *make_pass_omp_device_lower (gcc::context *ctxt); extern gimple_opt_pass *make_pass_object_sizes (gcc::context *ctxt); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c index d45326488cd..17cc9bd663e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c @@ -2,10 +2,10 @@ /* { dg-additional-options "-DVECTOR_LENGTH=" } */ /* { dg-additional-options "-fopenacc-dim=::128" } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include "pr85486.c" -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c index 33480a4ae68..5d05540ce46 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c @@ -2,10 +2,10 @@ /* { dg-additional-options "-DVECTOR_LENGTH=" } */ /* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include "pr85486.c" -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c index 0d98b82f993..f95f2ee3123 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c @@ -1,7 +1,7 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-DVECTOR_LENGTH=vector_length(128)" } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ /* Minimized from ref-1.C. */ @@ -54,5 +54,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c index dc92727d5be..7dc7459e5fe 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c @@ -4,7 +4,7 @@ { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } } */ -/* { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-fdump-tree-oaccloops" } */ /* { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'. */ @@ -36,9 +36,9 @@ static int fact_nohost(int n) return fact(n); } -/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target c } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'int fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && { ! offloading_enabled } } } } } - { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && offloading_enabled } } } } +/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccloops { target c } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'int fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccloops { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccloops { target { c++ && offloading_enabled } } } } TODO See PR101551 for 'offloading_enabled' differences. */ int main() diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c index 18d77cc5ecb..5158bb5eb89 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c @@ -1,5 +1,5 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -34,5 +34,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c index 8b5b2a4a92d..a3e44ebfbcb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c @@ -1,6 +1,6 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-fopenacc-dim=::128" } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -35,5 +35,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c index 59be37a7c27..a85400d09c5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c @@ -1,5 +1,5 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* We default to warp size 32 for the vector length, so the GOMP_OPENACC_DIM has no effect. */ /* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */ @@ -38,5 +38,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c index e5d1df09b8a..24c078f377c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c @@ -1,5 +1,5 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -36,5 +36,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c index e60f1c28db4..fcca9f593bb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c @@ -1,6 +1,6 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-fopenacc-dim=:2:128" } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -37,5 +37,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c index a1f67622f84..0807eab7eee 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c @@ -1,6 +1,6 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-set-target-env-var "GOMP_OPENACC_DIM" ":2:" } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -37,5 +37,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c index c419f6499b5..4a8c1bf549e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c @@ -1,5 +1,5 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -36,5 +36,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */ /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 index cd5bddc8685..b0537b8ff0b 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 @@ -5,7 +5,7 @@ ! With optimizations disabled, we currently don't expect that 'acc_on_device' "evaluates at compile time to a constant". ! { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } } -! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-fdump-tree-oaccloops" } program main use openacc @@ -58,6 +58,6 @@ function fact_nohost(x) result(res) res = fact(x) end function fact_nohost -! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } } -! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost_' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } } +! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } } +! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost_' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } } !TODO See PR101551 for 'offloading_enabled' differences. -- 2.30.2