From patchwork Wed Dec 15 15:54:21 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48957 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 938E8385802D for ; Wed, 15 Dec 2021 16:04:58 +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 CDB45385842C for ; Wed, 15 Dec 2021 15:55:53 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org CDB45385842C 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: bMv0WzH332lWgwmSSmY8dZBab5tSpPDUqxSK2VOxG1uYFnf9hnU4Ewreukk/OHPL9L9aLTL8mk BDaLL2KsY8sNTsgsDxu5CpqmsBsjvs2VJx+BAYxZv1gsPzuwexg15MjzQYyB9LZzk6seI1Ih6M jCJgPpc/zLGmLmHtmvy+Yon1kicrPhfVIpGCAECw9EfM/WZX6gRHLWSVnohHnOTMSa5zEtZjd3 MeyYiIsLv7JlbzNeZh0fdY+bd1f7yEsDQXuIBtJRs7df4YCkxBT+osU8jlly0lbekskLcXsyYY 95JrMPwUERRIyrkBOz/WyheG X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69736560" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:53 -0800 IronPort-SDR: gWy2gxaCFU4fDk1O2tZ/jIQrX17jWOn1rBaSFhBGJYWu3zFZXXa1GEodHVbILjhfBj9kG6lDHR APeaYp5cWt0p/uQBo6+FKuuQVBEL5lFA85FYay6uThGf5hZUZP9o7T+AiI5KzFZXka3lMW53Xu Cu1PmtaaKmU3nPO2rP1kvd/uKmrpDlzX3KK6+5865ntW37oybeaFN4tv6YQw6+sWYFaCy85sv5 LtnCj4M+SGKa1ryQZE8uuwzGUH9l3J3vVye8c7jzQVkYZHMgFaGKZ5rbMQYCWhHqHXNCyl/nyY lUA= From: Frederik Harwath To: Subject: [PATCH 14/40] openacc: Move pass_oacc_device_lower after pass_graphite Date: Wed, 15 Dec 2021 16:54:21 +0100 Message-ID: <20211215155447.19379-15-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com> References: <20211215155447.19379-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.6 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: rguenther@suse.de, thomas@codesourcery.com Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" The OpenACC device lowering pass must run after the Graphite pass to allow for the use of Graphite for automatic parallelization of kernels regions in the future. Experimentation has shown that it is best, performancewise, to run pass_oacc_device_lower together with the related passes pass_oacc_loop_designation and pass_oacc_gimple_workers early after pass_graphite in pass_tree_loop, at least if the other tree loop passes are not adjusted. In particular, to enable vectorization which is crucial for GCN offloading, device lowering should happen before pass_vectorize. To bring the loops contained in the offloading functions into the shape expected by the loop vectorizer, we have to make sure that some passes that previously were executed only once before pass_tree_loop are also executed on the offloading functions. To ensure the execution of pass_oacc_device_lower if pass_tree_loop does not execute (no loops, no optimizations), we introduce two further copies of the pass to the pipeline that run if there are no loops or if no optimization is performed. gcc/ChangeLog: * omp-general.c (oacc_get_fn_dim_size): Return 0 on missing "dims". * omp-oacc-neuter-broadcast.cc: Make pass_omp_oacc_neuter_broadcast clonable. * omp-offload.c (pass_oacc_loop_designation::clone): New member function. (pass_oacc_gimple_workers::clone): Likewise. (pass_oacc_gimple_device_lower::clone): Likewise. * passes.c (pass_data_no_loop_optimizations): New pass_data. (class pass_no_loop_optimizations): New pass. (make_pass_no_loop_optimizations): New function. * passes.def: Move pass_oacc_{loop_designation, gimple_workers, device_lower} into tree_loop, and add copies to pass_tree_no_loop and to new pass_no_loop_optimizations. Add copies of passes pass_ccp, pass_ipa_warn, pass_complete_unrolli, pass_backprop, pass_phiprop, pass_fix_loops after the OpenACC passes in pass_tree_loop. * tree-ssa-loop-ivcanon.c (pass_complete_unroll::clone): New member function. (pass_complete_unrolli::clone): Likewise. * tree-ssa-loop.c (pass_fix_loops::clone): Likewise. (pass_tree_loop_init::clone): Likewise. (pass_tree_loop_done::clone): Likewise. * tree-ssa-phiprop.c (pass_phiprop::clone): Likewise. * tree-pass.h (make_pass_oacc_only): New declaration. (make_pass_oacc_functions_only): New declaration. libgomp/ChangeLog: * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Adjust expected output to pass name changes due to the pass reordering and cloning. * 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. gcc/testsuite/ChangeLog: * gcc.dg/goacc/loop-processing-1.c: Adjust expected output to pass name changes due to the pass reordering and cloning. * c-c++-common/goacc/classify-kernels-unparallelized.c: Likewise. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/classify-parallel.c: Likewise. * c-c++-common/goacc/classify-routine.c: Likewise. * c-c++-common/goacc/routine-nohost-1.c: Likewise. * c-c++-common/unroll-1.c: Likewise. * c-c++-common/unroll-4.c: Likewise. * gcc.dg/tree-ssa/backprop-1.c: Likewise. * gcc.dg/tree-ssa/backprop-2.c: Likewise. * gcc.dg/tree-ssa/backprop-3.c: Likewise. * gcc.dg/tree-ssa/backprop-4.c: Likewise. * gcc.dg/tree-ssa/backprop-5.c: Likewise. * gcc.dg/tree-ssa/backprop-6.c: Likewise. * gcc.dg/tree-ssa/cunroll-1.c: Likewise. * gcc.dg/tree-ssa/cunroll-3.c: Likewise. * gcc.dg/tree-ssa/cunroll-9.c: Likewise. * gcc.dg/tree-ssa/ldist-17.c: Likewise. * gcc.dg/tree-ssa/loop-38.c: Likewise. * gcc.dg/tree-ssa/pr21463.c: Likewise. * gcc.dg/tree-ssa/pr45427.c: Likewise. * gcc.dg/tree-ssa/pr61743-1.c: Likewise. * gcc.dg/unroll-2.c: Likewise. * gcc.dg/unroll-3.c: Likewise. * gcc.dg/unroll-4.c: Likewise. * gcc.dg/unroll-5.c: Likewise. * gcc.dg/vect/vect-profile-1.c: Likewise. * gcc.dg/tree-ssa/loopclosedphi.c: Likewise. * gcc.dg/tree-ssa/pr59597.c: Likewise. * gcc.dg/vect/bb-slp-59.c: Likewise. * c-c++-common/goacc/device-lowering-debug-optimization.c: New test. * c-c++-common/goacc/device-lowering-no-loops.c: New test. * c-c++-common/goacc/device-lowering-no-optimization.c: New test. Co-Authored-By: Thomas Schwinge --- gcc/omp-general.c | 8 +- gcc/omp-oacc-neuter-broadcast.cc | 2 + gcc/omp-offload.c | 6 ++ gcc/passes.c | 42 ++++++++ gcc/passes.def | 44 ++++++++- .../goacc/classify-kernels-unparallelized.c | 8 +- .../c-c++-common/goacc/classify-kernels.c | 8 +- .../c-c++-common/goacc/classify-parallel.c | 8 +- .../c-c++-common/goacc/classify-routine.c | 22 ++--- .../device-lowering-debug-optimization.c | 29 ++++++ .../goacc/device-lowering-no-loops.c | 17 ++++ .../goacc/device-lowering-no-optimization.c | 30 ++++++ .../c-c++-common/goacc/routine-nohost-1.c | 6 +- gcc/testsuite/c-c++-common/unroll-1.c | 8 +- gcc/testsuite/c-c++-common/unroll-4.c | 4 +- .../gcc.dg/goacc/loop-processing-1.c | 5 +- gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c | 6 +- gcc/testsuite/gcc.dg/tree-ssa/backprop-2.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/backprop-3.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/backprop-4.c | 6 +- gcc/testsuite/gcc.dg/tree-ssa/backprop-5.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/backprop-6.c | 6 +- gcc/testsuite/gcc.dg/tree-ssa/cunroll-1.c | 6 +- gcc/testsuite/gcc.dg/tree-ssa/cunroll-3.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/cunroll-9.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/ldist-17.c | 2 +- gcc/testsuite/gcc.dg/tree-ssa/loop-38.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/loopclosedphi.c | 2 +- gcc/testsuite/gcc.dg/tree-ssa/pr21463.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/pr45427.c | 4 +- gcc/testsuite/gcc.dg/tree-ssa/pr59597.c | 2 +- gcc/testsuite/gcc.dg/tree-ssa/pr61743-1.c | 2 +- gcc/testsuite/gcc.dg/unroll-2.c | 2 +- gcc/testsuite/gcc.dg/unroll-3.c | 4 +- gcc/testsuite/gcc.dg/unroll-4.c | 4 +- gcc/testsuite/gcc.dg/unroll-5.c | 4 +- gcc/testsuite/gcc.dg/vect/bb-slp-59.c | 2 +- gcc/testsuite/gcc.dg/vect/vect-profile-1.c | 2 +- gcc/tree-pass.h | 2 + gcc/tree-ssa-loop-ivcanon.c | 2 + gcc/tree-ssa-loop.c | 99 +++++++++++++++++++ gcc/tree-ssa-phiprop.c | 2 + .../libgomp.oacc-c-c++-common/pr85486-2.c | 2 +- .../vector-length-128-1.c | 2 +- .../vector-length-128-2.c | 3 +- .../vector-length-128-3.c | 2 +- .../vector-length-128-4.c | 2 +- .../vector-length-128-5.c | 2 +- .../vector-length-128-6.c | 2 +- .../vector-length-128-7.c | 2 +- 50 files changed, 363 insertions(+), 88 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c create mode 100644 gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c create mode 100644 gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c -- 2.33.0 ----------------- 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 diff --git a/gcc/omp-general.c b/gcc/omp-general.c index 445275524134..27a1bc8092c8 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -2954,7 +2954,13 @@ oacc_get_fn_dim_size (tree fn, int axis) while (axis--) dims = TREE_CHAIN (dims); - int size = TREE_INT_CST_LOW (TREE_VALUE (dims)); + tree v = TREE_VALUE (dims); + /* TODO With 'pass_oacc_device_lower' moved "later", this is necessary to + avoid ICE for some OpenACC 'kernels' ("parloops") constructs. */ + if (v == NULL_TREE) + return 0; + + int size = TREE_INT_CST_LOW (v); return size; } diff --git a/gcc/omp-oacc-neuter-broadcast.cc b/gcc/omp-oacc-neuter-broadcast.cc index e43338f3abf2..94ecdc4d4e9a 100644 --- a/gcc/omp-oacc-neuter-broadcast.cc +++ b/gcc/omp-oacc-neuter-broadcast.cc @@ -1992,6 +1992,8 @@ public: return execute_omp_oacc_neuter_broadcast (); } + opt_pass * clone () { return new pass_omp_oacc_neuter_broadcast (m_ctxt); } + }; // class pass_omp_oacc_neuter_broadcast } // anon namespace diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 833f7ddea58f..e99aaac0e515 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -2444,6 +2444,8 @@ public: return execute_oacc_loop_designation (); } + opt_pass * clone () { return new pass_oacc_loop_designation (m_ctxt); } + }; // class pass_oacc_loop_designation const pass_data pass_data_oacc_device_lower = @@ -2467,12 +2469,16 @@ public: {} /* opt_pass methods: */ + /* TODO If this were gated on something like '!(fun->curr_properties & + PROP_gimple_oaccdevlow)', then we could easily have several instances + in the pass pipeline? */ virtual bool gate (function *) { return flag_openacc; }; virtual unsigned int execute (function *) { return execute_oacc_device_lower (); } + opt_pass * clone () { return new pass_oacc_device_lower (m_ctxt); } }; // class pass_oacc_device_lower diff --git a/gcc/passes.c b/gcc/passes.c index 64550b00b43c..4a1f4a4b5900 100644 --- a/gcc/passes.c +++ b/gcc/passes.c @@ -620,6 +620,48 @@ make_pass_all_optimizations_g (gcc::context *ctxt) namespace { +const pass_data pass_data_no_loop_optimizations = +{ + GIMPLE_PASS, /* type */ + "*no_loop_optimizations", /* name */ + OPTGROUP_NONE, /* optinfo_flags */ + TV_OPTIMIZE, /* tv_id */ + 0, /* properties_required */ + 0, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + 0, /* todo_flags_finish */ +}; + +/* This pass runs if loop optimizations are disabled + at the current optimization level. */ + +class pass_no_loop_optimizations : public gimple_opt_pass +{ +public: + pass_no_loop_optimizations (gcc::context *ctxt) + : gimple_opt_pass (pass_data_no_loop_optimizations, ctxt) + {} + + /* opt_pass methods: */ + virtual bool + gate (function *) + { + return !optimize || optimize_debug; + } + +}; // class pass_no_loop_optimizations + +} // anon namespace + +static gimple_opt_pass * +make_pass_no_loop_optimizations (gcc::context *ctxt) +{ + return new pass_no_loop_optimizations (ctxt); +} + +namespace { + const pass_data pass_data_rest_of_compilation = { RTL_PASS, /* type */ diff --git a/gcc/passes.def b/gcc/passes.def index 0f541454e7f1..5b9bb422d281 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -183,9 +183,6 @@ 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_omp_oacc_neuter_broadcast); - NEXT_PASS (pass_oacc_device_lower); NEXT_PASS (pass_omp_device_lower); NEXT_PASS (pass_omp_target_link); NEXT_PASS (pass_adjust_alignment); @@ -292,6 +289,35 @@ along with GCC; see the file COPYING3. If not see POP_INSERT_PASSES () NEXT_PASS (pass_parallelize_loops, false /* oacc_kernels_p */); NEXT_PASS (pass_expand_omp_ssa); + + /* Interrupt pass_tree_loop for OpenACC device lowering. */ + NEXT_PASS (pass_oacc_only); + PUSH_INSERT_PASSES_WITHIN (pass_oacc_only) + NEXT_PASS (pass_tree_loop_done); + NEXT_PASS (pass_oacc_loop_designation); + NEXT_PASS (pass_omp_oacc_neuter_broadcast); + NEXT_PASS (pass_oacc_device_lower); + + NEXT_PASS (pass_oacc_functions_only); + PUSH_INSERT_PASSES_WITHIN (pass_oacc_functions_only) + /* Repeat some passes on OpenACC functions after device lowering. */ + /* Lower complex instructions arising from OpenACC + reductions. */ + NEXT_PASS (pass_lower_complex); + /* Those passes are necessary here to allow the loop vectorizer to + work on the offloading functions which is important for AMD GCN + offloading. */ + NEXT_PASS (pass_ccp, true /* nonzero_p */); + NEXT_PASS (pass_complete_unrolli); + NEXT_PASS (pass_backprop); + NEXT_PASS (pass_phiprop); + NEXT_PASS (pass_fix_loops); + POP_INSERT_PASSES () + + /* Continue pass_tree_loop after OpenACC device lowering. */ + NEXT_PASS (pass_tree_loop_init); + POP_INSERT_PASSES () + NEXT_PASS (pass_ch_vect); NEXT_PASS (pass_if_conversion); /* pass_vectorize must immediately follow pass_if_conversion. @@ -311,15 +337,21 @@ along with GCC; see the file COPYING3. If not see NEXT_PASS (pass_loop_prefetch); /* Run IVOPTs after the last pass that uses data-reference analysis as that doesn't handle TARGET_MEM_REFs. */ + NEXT_PASS (pass_iv_optimize); NEXT_PASS (pass_lim); NEXT_PASS (pass_tree_loop_done); POP_INSERT_PASSES () + + /* Pass group that runs when pass_tree_loop is disabled or there are no loops in the function. */ NEXT_PASS (pass_tree_no_loop); PUSH_INSERT_PASSES_WITHIN (pass_tree_no_loop) NEXT_PASS (pass_slp_vectorize); + NEXT_PASS (pass_oacc_loop_designation); + NEXT_PASS (pass_omp_oacc_neuter_broadcast); + NEXT_PASS (pass_oacc_device_lower); POP_INSERT_PASSES () NEXT_PASS (pass_simduid_cleanup); NEXT_PASS (pass_lower_vector_ssa); @@ -397,6 +429,12 @@ along with GCC; see the file COPYING3. If not see NEXT_PASS (pass_local_pure_const); NEXT_PASS (pass_modref); POP_INSERT_PASSES () + NEXT_PASS (pass_no_loop_optimizations); + PUSH_INSERT_PASSES_WITHIN (pass_no_loop_optimizations) + NEXT_PASS (pass_oacc_loop_designation); + NEXT_PASS (pass_omp_oacc_neuter_broadcast); + NEXT_PASS (pass_oacc_device_lower); + POP_INSERT_PASSES () NEXT_PASS (pass_tm_init); PUSH_INSERT_PASSES_WITHIN (pass_tm_init) NEXT_PASS (pass_tm_mark); 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 e391184f403d..338676aa20ff 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -6,7 +6,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-oaccloops" } */ + { dg-additional-options "-fdump-tree-oaccloops1" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -39,6 +39,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 "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" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops1" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 779e2b0a24db..37e2a57455d1 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -6,7 +6,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-oaccloops" } */ + { dg-additional-options "-fdump-tree-oaccloops1" } */ /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting aspects of that functionality. */ @@ -35,6 +35,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 "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" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops1" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 9056aa69dad6..82b70ae280cd 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-oaccloops" } */ + { dg-additional-options "-fdump-tree-oaccloops1" } */ /* { 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 "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" } } */ + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops1" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index f7f0454009bf..cd539370dbbf 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-oaccloops" } */ + { dg-additional-options "-fdump-tree-oaccloops1" } */ /* { 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 "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 } } } } + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccloops1" { target c } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops1" { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops1" { target { c++ && offloading_enabled } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccloops1" { target c } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccloops1" { target { c++ && { ! offloading_enabled } } } } } + { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccloops1" { target { c++ && offloading_enabled } } } } TODO See PR101551 for 'offloading_enabled' differences. - { 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" } } */ + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } } + { 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 "oaccloops1" } } + { dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 1 "oaccloops1" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c b/gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c new file mode 100644 index 000000000000..5bf37cc61580 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c @@ -0,0 +1,29 @@ +/* Verify that OpenACC device lowering executes with "-Og". The actual logic in + the test function does not matter. */ + +/* { dg-additional-options "-Og -fdump-tree-oaccdevlow" } */ + +int main() +{ + int i, j; + int ina[1024], out[1024], acc; + + for (j = 0; j < 32; j++) + for (i = 0; i < 32; i++) + ina[j * 32 + i] = (i == j) ? 2 : 0; + + acc = 0; +#pragma acc parallel loop copy(acc, ina, out) + for (j = 0; j < 32; j++) + { +#pragma acc loop reduction(+:acc) + for (i = 0; i < 32; i++) + acc += ina[i]; + + out[j] = acc; + } + + return 0; +} + +/* { dg-final { scan-tree-dump ".omp_fn" "oaccdevlow3" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c new file mode 100644 index 000000000000..193b5620de1d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c @@ -0,0 +1,17 @@ +/* Verify that OpenACC device lowering executes even if there are no OpenACC + loops. */ + +/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow" } */ + +int main() +{ + int x; +#pragma acc parallel copy(x) + { + asm volatile(""); + } + + return 0; +} + +/* { dg-final { scan-tree-dump ".omp_fn" "oaccdevlow2" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c new file mode 100644 index 000000000000..69e2b22d73ba --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c @@ -0,0 +1,30 @@ +/* Verify that OpenACC device lowering executes with "-O0". The actual + logic in the test function does not matter. */ + +/* { dg-additional-options "-O0 -fdump-tree-oaccdevlow" } */ + +int main() +{ + + int i, j; + int ina[1024], out[1024], acc; + + for (j = 0; j < 32; j++) + for (i = 0; i < 32; i++) + ina[j * 32 + i] = (i == j) ? 2 : 0; + + acc = 0; +#pragma acc parallel loop copy(acc, ina, out) + for (j = 0; j < 32; j++) + { +#pragma acc loop reduction(+:acc) + for (i = 0; i < 32; i++) + acc += ina[i]; + + out[j] = acc; + } + + return 0; +} + +/* { dg-final { scan-tree-dump ".omp_fn" "oaccdevlow3" } } */ 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 59ebb2bc5a9f..4f9a3a333570 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c @@ -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 oaccloops } } */ +/* { 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 oaccloops } } */ +/* { 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 oaccloops } } */ +/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 "oaccloops*" } } */ diff --git a/gcc/testsuite/c-c++-common/unroll-1.c b/gcc/testsuite/c-c++-common/unroll-1.c index fe7f4f31912c..8e57a44be231 100644 --- a/gcc/testsuite/c-c++-common/unroll-1.c +++ b/gcc/testsuite/c-c++-common/unroll-1.c @@ -1,5 +1,5 @@ -/* { dg-do compile } */ -/* { dg-options "-O2 -fdump-tree-cunrolli-details -fdump-rtl-loop2_unroll-details" } */ +/* { dg-do compile } * +/* { dg-options "-O2 -fdump-tree-cunrolli1-details -fdump-rtl-loop2_unroll-details" } */ extern void bar (int); @@ -10,12 +10,12 @@ void test (void) #pragma GCC unroll 8 for (unsigned long i = 1; i <= 8; ++i) bar(i); - /* { dg-final { scan-tree-dump "11:.*: loop with 8 iterations completely unrolled" "cunrolli" } } */ + /* { dg-final { scan-tree-dump "11:.*: loop with 8 iterations completely unrolled" "cunrolli1" } } */ #pragma GCC unroll 8 for (unsigned long i = 1; i <= 7; ++i) bar(i); - /* { dg-final { scan-tree-dump "16:.*: loop with 7 iterations completely unrolled" "cunrolli" } } */ + /* { dg-final { scan-tree-dump "16:.*: loop with 7 iterations completely unrolled" "cunrolli1" } } */ #pragma GCC unroll 8 for (unsigned long i = 1; i <= 15; ++i) diff --git a/gcc/testsuite/c-c++-common/unroll-4.c b/gcc/testsuite/c-c++-common/unroll-4.c index 1c1988174ba7..fe7f9e10626e 100644 --- a/gcc/testsuite/c-c++-common/unroll-4.c +++ b/gcc/testsuite/c-c++-common/unroll-4.c @@ -1,5 +1,5 @@ /* { dg-do compile } */ -/* { dg-options "-O2 -funroll-all-loops -fdump-rtl-loop2_unroll-details -fdump-tree-cunrolli-details" } */ +/* { dg-options "-O2 -funroll-all-loops -fdump-rtl-loop2_unroll-details -fdump-tree-cunrolli1-details" } */ extern void bar (int); @@ -17,6 +17,6 @@ void test (void) for (unsigned long i = 1; i <= j; ++i) bar(i); - /* { dg-final { scan-tree-dump "Not unrolling loop .: user didn't want it unrolled completely" "cunrolli" } } */ + /* { dg-final { scan-tree-dump "Not unrolling loop .: user didn't want it unrolled completely" "cunrolli1" } } */ /* { dg-final { scan-rtl-dump-times "Not unrolling loop, user didn't want it unrolled" 2 "loop2_unroll" } } */ } diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c index 78b9aed89beb..c191125b7951 100644 --- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c +++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -1,5 +1,4 @@ -/* Make sure that OpenACC loop processing happens. */ -/* { dg-additional-options "-O2 -fdump-tree-oaccloops" } */ +/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow*" } */ extern int place (); @@ -15,4 +14,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\);} "oaccloops" } } */ +/* { 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/gcc.dg/tree-ssa/backprop-1.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c index 302fdb570b63..b6b11bf30afa 100644 --- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c +++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c @@ -1,5 +1,5 @@ /* { dg-do compile } */ -/* { dg-options "-O -g -fdump-tree-backprop-details" } */ +/* { dg-options "-O -g -fdump-tree-backprop1-details" } */ /* Test a simple case of non-looping code in which both uses ignore the sign and both definitions are sign ops. */ @@ -18,5 +18,5 @@ TEST_FUNCTION (float, f) TEST_FUNCTION (double, ) TEST_FUNCTION (long double, l) -/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 3 "backprop" } } */ -/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = ABS_EXPR = 2; } virtual unsigned int execute (function *); + opt_pass * clone () { return new pass_complete_unrolli (m_ctxt); } }; // class pass_complete_unrolli diff --git a/gcc/tree-ssa-loop.c b/gcc/tree-ssa-loop.c index 1bbf2f1fb2c8..8d5572033f7b 100644 --- a/gcc/tree-ssa-loop.c +++ b/gcc/tree-ssa-loop.c @@ -70,6 +70,8 @@ public: virtual bool gate (function *) { return flag_tree_loop_optimize; } virtual unsigned int execute (function *fn); + + opt_pass * clone () { return new pass_fix_loops (m_ctxt); } }; // class pass_fix_loops unsigned int @@ -136,6 +138,8 @@ public: /* opt_pass methods: */ virtual bool gate (function *fn) { return gate_loop (fn); } + + opt_pass * clone () { return new pass_tree_loop (m_ctxt); } }; // class pass_tree_loop } // anon namespace @@ -200,6 +204,97 @@ make_pass_oacc_kernels (gcc::context *ctxt) { return new pass_oacc_kernels (ctxt); } +/* A superpass that runs its subpasses on OpenACC functions only. */ + +namespace { + +const pass_data pass_data_oacc_functions_only = +{ + GIMPLE_PASS, /* type */ + "*oacc_fns_only", /* name */ + OPTGROUP_LOOP, /* optinfo_flags */ + TV_TREE_LOOP, /* tv_id */ + 0, /* properties_required */ + 0, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + 0, /* todo_flags_finish */ +}; + +class pass_oacc_functions_only: public gimple_opt_pass +{ +public: + pass_oacc_functions_only (gcc::context *ctxt) + : gimple_opt_pass (pass_data_oacc_functions_only, ctxt) + {} + + /* opt_pass methods: */ + virtual bool gate (function *fn) { + if (!flag_openacc) + return false; + + if (!oacc_get_fn_attrib (fn->decl)) + return false; + + return true; + } + +}; // class pass_oacc_functions_only + +} // anon namespace + +gimple_opt_pass * +make_pass_oacc_functions_only (gcc::context *ctxt) +{ + return new pass_oacc_functions_only (ctxt); +} + +/* A superpass that runs its subpasses only if compiling for OpenACC. */ + +namespace { + +const pass_data pass_data_oacc_only = +{ + GIMPLE_PASS, /* type */ + "*oacc_only", /* name */ + OPTGROUP_LOOP, /* optinfo_flags */ + TV_TREE_LOOP, /* tv_id */ + 0, /* properties_required */ + 0, /* properties_provided */ + 0, /* properties_destroyed */ + 0, /* todo_flags_start */ + 0, /* todo_flags_finish */ +}; + +class pass_oacc_only: public gimple_opt_pass +{ +public: + pass_oacc_only (gcc::context *ctxt) + : gimple_opt_pass (pass_data_oacc_only, ctxt) + {} + + /* opt_pass methods: */ + virtual bool gate (function *fn) { + if (!flag_openacc) + return false; + + if (!oacc_get_fn_attrib (fn->decl)) + return false; + + return true; + } + +}; // class pass_oacc_only + +} // anon namespace + +gimple_opt_pass * +make_pass_oacc_only (gcc::context *ctxt) +{ + return new pass_oacc_only (ctxt); +} + + /* The ipa oacc superpass. */ @@ -343,6 +438,8 @@ public: /* opt_pass methods: */ virtual unsigned int execute (function *); + opt_pass * clone () { return new pass_tree_loop_init (m_ctxt); } + }; // class pass_tree_loop_init unsigned int @@ -556,6 +653,8 @@ public: /* opt_pass methods: */ virtual unsigned int execute (function *) { return tree_ssa_loop_done (); } + opt_pass * clone () { return new pass_tree_loop_done (m_ctxt); } + }; // class pass_tree_loop_done } // anon namespace diff --git a/gcc/tree-ssa-phiprop.c b/gcc/tree-ssa-phiprop.c index 78b0461c839d..f138f766286b 100644 --- a/gcc/tree-ssa-phiprop.c +++ b/gcc/tree-ssa-phiprop.c @@ -479,6 +479,8 @@ public: virtual bool gate (function *) { return flag_tree_phiprop; } virtual unsigned int execute (function *); + opt_pass * clone () { return new pass_phiprop (m_ctxt); } + }; // class pass_phiprop unsigned int 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 17cc9bd663e5..4438f6c24fed 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c @@ -7,5 +7,5 @@ #include "pr85486.c" -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops1" } } */ /* { 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-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c index 5158bb5eb89e..c0a29c7556f9 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 @@ -34,5 +34,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */ +/* { 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 a3e44ebfbcb4..326f6d8dc31a 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,7 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-fopenacc-dim=::128" } */ /* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */ +>>>>>>> adfd567486a0 (Move pass_oacc_device_lower after pass_graphite) /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ #include @@ -35,5 +36,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */ +/* { 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 a85400d09c50..efc9297acdee 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 @@ -38,5 +38,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */ +/* { 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 24c078f377c3..1c83ec0cc18d 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 @@ -36,5 +36,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */ +/* { 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 fcca9f593bb2..f2391dca7272 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 @@ -37,5 +37,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */ +/* { 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 0807eab7eee4..8ddaaf592cc1 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 @@ -37,5 +37,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */ +/* { 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 4a8c1bf549e9..97abbfc20986 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 @@ -36,5 +36,5 @@ main (void) return 0; } -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */ +/* { 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" } */