From patchwork Wed Dec 15 15:54:08 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48944 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 81C9C385842F for ; Wed, 15 Dec 2021 15:56:37 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 3775B3858C27 for ; Wed, 15 Dec 2021 15:55:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 3775B3858C27 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: mbhOgLMGLXef7wf8fXjfOWOUB2F/2rmiEguK7ET9U1Nc0YJz/gxiEafLTmgRraNQtNrMGUOpw+ JU2NdsB8H8hEgOfTVKQ7+2MvjX+YAEBinWe8unK9dPXzA6CiGSvu1mKmgNqbVZzsgTXfsyru8Z xSWzUITAN7/5vxxAsqIctVv/4encFby+NMsj2NXyqX8oTxe+8Mmg8SQ7kphrHlPqjX+emZfl9l jM/XCE2YGaSDDPlDED+fngQJS4GQrChFFV+uHZry1XE4QC4GRIVP7H6/4RzNESl+yPUov0SQNG Qy2awrvHukWGdMnBXh7/q0Ld X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69584563" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:01 -0800 IronPort-SDR: xca62t3aRMYV59bLcbrlIXvGu7BJFcP6pFmcf21Pf+qaCmsNvDd61/217up6VO+bLg1+9O7wm9 HWWkO7wJ7k20yaMUZr8dUnDUjodYhyOGaFXE72/4oI4GTvk1lQZzMSDLCDxrRTzk99j3/QqNoJ QRC0v3QfDUHfpQEikJvWVp2m2AfHtfqDKshBdCBd5dmVbb/HLwXhHoshYDIFO7ZkgkJTUx/qwm cI9rtrXRF/zS57/0Ppohe0FO1ki0uQrrCKEgdBfbJZIC1aMWQiIMUsMzwFKLlhdcjMNv+8yLxg sjA= From: Frederik Harwath To: Subject: [PATCH 01/40] Kernels loops annotation: C and C++. Date: Wed, 15 Dec 2021 16:54:08 +0100 Message-ID: <20211215155447.19379-2-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-09.mgc.mentorg.com (139.181.222.9) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-9.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPAM_BODY, 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: Sandra Loosemore , thomas@codesourcery.com, joseph@codesourcery.com, nathan@acm.org Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" From: Sandra Loosemore This patch detects loops in kernels regions that are candidates for parallelization, and adds "#pragma acc loop auto" annotations to them. This annotation is controlled by the -fopenacc-kernels-annotate-loops option, which is enabled by default. -Wopenacc-kernels-annotate-loops can be used to produce diagnostics about loops that cannot be annotated. gcc/c-family/ * c-common.h (c_oacc_annotate_loops_in_kernels_regions): Declare. * c-omp.c: Include tree-iterator.h (enum annotation_state): New. (struct annotation_info): New. (do_not_annotate_loop): New. (do_not_annotate_loop_nest): New. (annotation_error): New. (c_finish_omp_for_internal): Split from c_finish_omp_for. Use annotation_error function. Code refactoring to avoid destructive changes that cannot be undone in case of error. (is_local_var): New. (lang_specific_unwrap_initializer): New. (annotate_for_loop): New. (check_and_annotate_for_loop): New. (annotate_loops_in_kernels_regions): New. (c_oacc_annotate_loops_in_kernels_regions): New. * c.opt (Wopenacc-kernels-annotate-loops): New. (fopenacc-kernels-annotate-loops): New. gcc/c/ * c-decl.c (c_unwrap_for_init): New. (finish_function): Call c_oacc_annotate_loops_in_kernels_regions. gcc/cp/ * decl.c (cp_unwrap_for_init): New. (finish_function): Call c_oacc_annotate_loops_in_kernels_regions. gcc/ * doc/invoke.texi (Option Summary): Add entries for -Wopenacc-kernels-annotate-loops and -fno-openacc-kernels-annotate-loops. (Warning Options): Document -Wopenacc-kernels-annotate-loops. (Optimization Options): Document -fno-openacc-kernels-annotate-loops. gcc/testsuite/ * c-c++-common/goacc/classify-kernels-unparallelized.c: Add -fno-openacc-kernels-annotate-loops option. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/kernels-counter-var-redundant-load.c: Likewise. * c-c++-common/goacc/kernels-counter-vars-function-scope.c: Likewise. * c-c++-common/goacc/kernels-double-reduction.c: Likewise. * c-c++-common/goacc/kernels-double-reduction-n.c: Likewise. * c-c++-common/goacc/kernels-loop-2.c: Likewise. * c-c++-common/goacc/kernels-loop-3.c: Likewise. * c-c++-common/goacc/kernels-loop-data-2.c: Likewise. * c-c++-common/goacc/kernels-loop-data-enter-exit-2.c: Likewise. * c-c++-common/goacc/kernels-loop-data-enter-exit.c: Likewise. * c-c++-common/goacc/kernels-loop-data-update.c: Likewise. * c-c++-common/goacc/kernels-loop-data.c: Likewise. * c-c++-common/goacc/kernels-loop-g.c: Likewise. * c-c++-common/goacc/kernels-loop-mod-not-zero.c: Likewise. * c-c++-common/goacc/kernels-loop-n.c: Likewise. * c-c++-common/goacc/kernels-loop-nest.c: Likewise. * c-c++-common/goacc/kernels-loop.c: Likewise. * c-c++-common/goacc/kernels-one-counter-var.c: Likewise. * c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c: Likewise. * c-c++-common/goacc/kernels-reduction.c: Likewise. * c-c++-common/goacc/kernels-loop-annotation-1.c: New. * c-c++-common/goacc/kernels-loop-annotation-2.c: New. * c-c++-common/goacc/kernels-loop-annotation-3.c: New. * c-c++-common/goacc/kernels-loop-annotation-4.c: New. * c-c++-common/goacc/kernels-loop-annotation-5.c: New. * c-c++-common/goacc/kernels-loop-annotation-6.c: New. * c-c++-common/goacc/kernels-loop-annotation-7.c: New. * c-c++-common/goacc/kernels-loop-annotation-8.c: New. * c-c++-common/goacc/kernels-loop-annotation-9.c: New. * c-c++-common/goacc/kernels-loop-annotation-10.c: New. * c-c++-common/goacc/kernels-loop-annotation-11.c: New. * c-c++-common/goacc/kernels-loop-annotation-12.c: New. * c-c++-common/goacc/kernels-loop-annotation-13.c: New. * c-c++-common/goacc/kernels-loop-annotation-14.c: New. * c-c++-common/goacc/kernels-loop-annotation-15.c: New. * c-c++-common/goacc/kernels-loop-annotation-16.c: New. * c-c++-common/goacc/kernels-loop-annotation-17.c: New. --- gcc/c-family/c-common.h | 1 + gcc/c-family/c-omp.c | 799 ++++++++++++++++-- gcc/c-family/c.opt | 8 + gcc/c/c-decl.c | 28 + gcc/cp/decl.c | 44 + gcc/doc/invoke.texi | 32 +- .../goacc/classify-kernels-unparallelized.c | 1 + .../c-c++-common/goacc/classify-kernels.c | 3 +- .../kernels-counter-var-redundant-load.c | 1 + .../kernels-counter-vars-function-scope.c | 1 + .../goacc/kernels-double-reduction-n.c | 1 + .../goacc/kernels-double-reduction.c | 1 + .../c-c++-common/goacc/kernels-loop-2.c | 1 + .../c-c++-common/goacc/kernels-loop-3.c | 1 + .../goacc/kernels-loop-annotation-1.c | 26 + .../goacc/kernels-loop-annotation-10.c | 32 + .../goacc/kernels-loop-annotation-11.c | 27 + .../goacc/kernels-loop-annotation-12.c | 28 + .../goacc/kernels-loop-annotation-13.c | 27 + .../goacc/kernels-loop-annotation-14.c | 22 + .../goacc/kernels-loop-annotation-15.c | 22 + .../goacc/kernels-loop-annotation-16.c | 26 + .../goacc/kernels-loop-annotation-17.c | 26 + .../goacc/kernels-loop-annotation-2.c | 21 + .../goacc/kernels-loop-annotation-3.c | 24 + .../goacc/kernels-loop-annotation-4.c | 34 + .../goacc/kernels-loop-annotation-5.c | 27 + .../goacc/kernels-loop-annotation-6.c | 27 + .../goacc/kernels-loop-annotation-7.c | 26 + .../goacc/kernels-loop-annotation-8.c | 27 + .../goacc/kernels-loop-annotation-9.c | 26 + .../c-c++-common/goacc/kernels-loop-data-2.c | 1 + .../goacc/kernels-loop-data-enter-exit-2.c | 1 + .../goacc/kernels-loop-data-enter-exit.c | 1 + .../goacc/kernels-loop-data-update.c | 1 + .../c-c++-common/goacc/kernels-loop-data.c | 1 + .../c-c++-common/goacc/kernels-loop-g.c | 1 + .../goacc/kernels-loop-mod-not-zero.c | 1 + .../c-c++-common/goacc/kernels-loop-n.c | 1 + .../c-c++-common/goacc/kernels-loop-nest.c | 1 + .../c-c++-common/goacc/kernels-loop.c | 1 + .../goacc/kernels-one-counter-var.c | 1 + .../kernels-parallel-loop-data-enter-exit.c | 1 + .../c-c++-common/goacc/kernels-reduction.c | 1 + 44 files changed, 1322 insertions(+), 61 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-10.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-11.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-12.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-13.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-14.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-15.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-16.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-17.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-3.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-4.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-5.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-6.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-7.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-8.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-9.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/c-family/c-common.h b/gcc/c-family/c-common.h index f60714e34160..f8b414401a5d 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1247,6 +1247,7 @@ extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree); extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree); extern tree c_omp_check_context_selector (location_t, tree); extern void c_omp_mark_declare_variant (location_t, tree, tree); +extern void c_oacc_annotate_loops_in_kernels_regions (tree, tree (*) (tree)); extern const char *c_omp_map_clause_name (tree, bool); extern void c_omp_adjust_map_clauses (tree, bool); diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index fad060670b65..fad50da8fbc4 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -37,7 +37,7 @@ along with GCC; see the file COPYING3. If not see #include "langhooks.h" #include "bitmap.h" #include "gimple-fold.h" - +#include "tree-iterator.h" /* Complete a #pragma oacc wait construct. LOC is the location of the #pragma. */ @@ -918,6 +918,110 @@ c_omp_for_incr_canonicalize_ptr (location_t loc, tree decl, tree incr) return incr; } +/* State of annotation traversal for FOR loops in kernels regions, + used to control processing and diagnostic messages that are deferred until + the entire loop has been scanned. */ +enum annotation_state { + as_outer, + as_in_kernels_region, + as_in_kernels_loop, + /* The remaining state values represent conversion failures caught + while in as_in_kernels_loop state. To test whether the traversal is + in the body of a kernels loop, use (state >= as_in_kernels_loop). */ + as_invalid_variable_type, + as_missing_initializer, + as_invalid_initializer, + as_missing_predicate, + as_invalid_predicate, + as_missing_increment, + as_invalid_increment, + as_explicit_annotation, + as_invalid_control_flow, + as_invalid_break, + as_invalid_return, + as_invalid_call, + as_invalid_modification +}; + +/* Structure used to hold state for automatic annotation of FOR loops + in kernels regions. LOOP is the nearest enclosing loop, or + NULL_TREE if outside of a loop context. VARS is a tree_list + containing the variables controlling LOOP's termination (the + induction variable and a possible limit variable). STATE keeps + track of whether loop satisfies all criteria making it legal to + parallelize. Otherwise, REASON is a statement that blocks + automatic parallelization, such as an unstructured jump or an + assignment to a variable in VARS, used for printing diagnostics. + + These structures are chained through NEXT, which points to the + next-closest enclosing loop's or the kernels region's annotation info, if + any. */ + +struct annotation_info +{ + tree loop; + tree vars; + bool break_ok; + enum annotation_state state; + tree reason; + struct annotation_info *next; +}; + +/* Mark the current loop's INFO as not OK to annotate, recording STATE + and REASON for producing diagnostics later. */ + +static void +do_not_annotate_loop (struct annotation_info *info, + enum annotation_state state, tree reason) +{ + if (info->state == as_in_kernels_loop) + { + info->state = state; + info->reason = reason; + } +} + +/* Mark the current loop identified by INFO and all of its ancestors (i.e., + enclosing loops) as not OK to annotate. Arguments are the same as + for do_not_annotate_loop. */ + +static void +do_not_annotate_loop_nest (struct annotation_info *info, + enum annotation_state state, tree reason) +{ + while (info != NULL) + { + do_not_annotate_loop (info, state, reason); + info = info->next; + } +} + +/* If INFO is non-null, call do_not_annotate_loop with STATE and REASON + to record info for diagnosing an error later. Otherwise emit an error now + at ELOCUS with message MSG and the optional arguments. */ + +static void annotation_error (struct annotation_info *, + enum annotation_state, tree, location_t, + const char *, ...) ATTRIBUTE_GCC_DIAG(5,6); +static +void annotation_error (struct annotation_info *info, + enum annotation_state state, + tree reason, + location_t elocus, + const char *msg, ...) +{ + if (info) + do_not_annotate_loop (info, state, reason); + else + { + auto_diagnostic_group d; + va_list ap; + va_start (ap, msg); + emit_diagnostic_valist (DK_ERROR, elocus, -1, msg, &ap); + va_end (ap); + } +} + /* Validate and generate OMP_FOR. DECLV is a vector of iteration variables, for each collapsed loop. @@ -927,12 +1031,19 @@ c_omp_for_incr_canonicalize_ptr (location_t loc, tree decl, tree incr) INITV, CONDV and INCRV are vectors containing initialization expressions, controlling predicates and increment expressions. BODY is the body of the loop and PRE_BODY statements that go before - the loop. */ + the loop. FINAL_P is true if not inside a C++ template. -tree -c_finish_omp_for (location_t locus, enum tree_code code, tree declv, - tree orig_declv, tree initv, tree condv, tree incrv, - tree body, tree pre_body, bool final_p) + INFO is null if called to parse an explicitly-annotated OMP for + loop, otherwise it holds state information for automatically + annotating a regular FOR loop in a kernels region. In the former case, + malformed loops are hard errors; otherwise we just record the annotation + failure in INFO. */ + +static tree +c_finish_omp_for_internal (location_t locus, enum tree_code code, tree declv, + tree orig_declv, tree initv, tree condv, tree incrv, + tree body, tree pre_body, bool final_p, + struct annotation_info *info) { location_t elocus; bool fail = false; @@ -956,12 +1067,14 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, if (!INTEGRAL_TYPE_P (TREE_TYPE (decl)) && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE) { - error_at (elocus, "invalid type for iteration variable %qE", decl); + annotation_error (info, as_invalid_variable_type, decl, elocus, + "invalid type for iteration variable %qE", decl); fail = true; } else if (TYPE_ATOMIC (TREE_TYPE (decl))) { - error_at (elocus, "%<_Atomic%> iteration variable %qE", decl); + annotation_error (info, as_invalid_variable_type, decl, elocus, + "%<_Atomic%> iteration variable %qE", decl); fail = true; /* _Atomic iterator confuses stuff too much, so we risk ICE trying to diagnose it further. */ @@ -977,7 +1090,8 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, init = DECL_INITIAL (decl); if (init == NULL) { - error_at (elocus, "%qE is not initialized", decl); + annotation_error (info, as_missing_initializer, decl, elocus, + "%qE is not initialized", decl); init = integer_zero_node; fail = true; } @@ -998,7 +1112,8 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, if (cond == NULL_TREE) { - error_at (elocus, "missing controlling predicate"); + annotation_error (info, as_missing_predicate, NULL_TREE, elocus, + "missing controlling predicate"); fail = true; } else @@ -1014,12 +1129,14 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, if (EXPR_HAS_LOCATION (cond)) elocus = EXPR_LOCATION (cond); - if (TREE_CODE (cond) == LT_EXPR - || TREE_CODE (cond) == LE_EXPR - || TREE_CODE (cond) == GT_EXPR - || TREE_CODE (cond) == GE_EXPR - || TREE_CODE (cond) == NE_EXPR - || TREE_CODE (cond) == EQ_EXPR) + enum tree_code condcode = TREE_CODE (cond); + + if (condcode == LT_EXPR + || condcode == LE_EXPR + || condcode == GT_EXPR + || condcode == GE_EXPR + || condcode == NE_EXPR + || condcode == EQ_EXPR) { tree op0 = TREE_OPERAND (cond, 0); tree op1 = TREE_OPERAND (cond, 1); @@ -1039,79 +1156,88 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, if (TREE_CODE (op0) == NOP_EXPR && decl == TREE_OPERAND (op0, 0)) { - TREE_OPERAND (cond, 0) = TREE_OPERAND (op0, 0); - TREE_OPERAND (cond, 1) - = fold_build1_loc (elocus, NOP_EXPR, TREE_TYPE (decl), - TREE_OPERAND (cond, 1)); + op0 = TREE_OPERAND (op0, 0); + op1 = fold_build1_loc (elocus, NOP_EXPR, TREE_TYPE (decl), + op1); } else if (TREE_CODE (op1) == NOP_EXPR && decl == TREE_OPERAND (op1, 0)) { - TREE_OPERAND (cond, 1) = TREE_OPERAND (op1, 0); - TREE_OPERAND (cond, 0) - = fold_build1_loc (elocus, NOP_EXPR, TREE_TYPE (decl), - TREE_OPERAND (cond, 0)); + op1 = TREE_OPERAND (op1, 0); + op0 = fold_build1_loc (elocus, NOP_EXPR, TREE_TYPE (decl), + op0); } - if (decl == TREE_OPERAND (cond, 0)) + if (decl == op0) cond_ok = true; - else if (decl == TREE_OPERAND (cond, 1)) + else if (decl == op1) { - TREE_SET_CODE (cond, - swap_tree_comparison (TREE_CODE (cond))); - TREE_OPERAND (cond, 1) = TREE_OPERAND (cond, 0); - TREE_OPERAND (cond, 0) = decl; + condcode = swap_tree_comparison (condcode); + op1 = op0; + op0 = decl; cond_ok = true; } - if (TREE_CODE (cond) == NE_EXPR - || TREE_CODE (cond) == EQ_EXPR) + if (condcode == NE_EXPR || condcode == EQ_EXPR) { if (!INTEGRAL_TYPE_P (TREE_TYPE (decl))) { - if (code == OACC_LOOP || TREE_CODE (cond) == EQ_EXPR) + if (code == OACC_LOOP || condcode == EQ_EXPR) cond_ok = false; } - else if (operand_equal_p (TREE_OPERAND (cond, 1), + else if (operand_equal_p (op1, TYPE_MIN_VALUE (TREE_TYPE (decl)), 0)) - TREE_SET_CODE (cond, TREE_CODE (cond) == NE_EXPR - ? GT_EXPR : LE_EXPR); - else if (operand_equal_p (TREE_OPERAND (cond, 1), + condcode = (condcode == NE_EXPR ? GT_EXPR : LE_EXPR); + else if (operand_equal_p (op1, TYPE_MAX_VALUE (TREE_TYPE (decl)), 0)) - TREE_SET_CODE (cond, TREE_CODE (cond) == NE_EXPR - ? LT_EXPR : GE_EXPR); - else if (code == OACC_LOOP || TREE_CODE (cond) == EQ_EXPR) + condcode = (condcode == NE_EXPR ? LT_EXPR : GE_EXPR); + else if (code == OACC_LOOP || condcode == EQ_EXPR) cond_ok = false; } - if (cond_ok && TREE_VEC_ELT (condv, i) != cond) + if (cond_ok) { - tree ce = NULL_TREE, *pce = &ce; - tree type = TREE_TYPE (TREE_OPERAND (cond, 1)); - for (tree c = TREE_VEC_ELT (condv, i); c != cond; - c = TREE_OPERAND (c, 1)) + /* We postponed destructive changes to canonicalize + cond until we're sure it is OK. In the !error_p + case where we are trying to transform a regular FOR_STMT + to OMP_FOR, we don't want to destroy the original + condition if we aren't going to be able to do the + transformation anyway. */ + TREE_SET_CODE (cond, condcode); + TREE_OPERAND (cond, 0) = op0; + TREE_OPERAND (cond, 1) = op1; + + if (TREE_VEC_ELT (condv, i) != cond) { - *pce = build2 (COMPOUND_EXPR, type, TREE_OPERAND (c, 0), - TREE_OPERAND (cond, 1)); - pce = &TREE_OPERAND (*pce, 1); + tree ce = NULL_TREE, *pce = &ce; + tree type = TREE_TYPE (op1); + for (tree c = TREE_VEC_ELT (condv, i); c != cond; + c = TREE_OPERAND (c, 1)) + { + *pce = build2 (COMPOUND_EXPR, type, + TREE_OPERAND (c, 0), op1); + pce = &TREE_OPERAND (*pce, 1); + } + op1 = ce; + TREE_VEC_ELT (condv, i) = cond; } - TREE_OPERAND (cond, 1) = ce; - TREE_VEC_ELT (condv, i) = cond; } } if (!cond_ok) { - error_at (elocus, "invalid controlling predicate"); + annotation_error (info, as_invalid_predicate, cond, elocus, + "invalid controlling predicate"); fail = true; } } if (incr == NULL_TREE) { - error_at (elocus, "missing increment expression"); + annotation_error (info, as_missing_increment, NULL_TREE, elocus, + "missing increment expression"); fail = true; } else @@ -1210,9 +1336,11 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, if (i == NULL_TREE || !operand_equal_p (unit, i, 0)) { - error_at (elocus, - "increment is not constant 1 or " - "-1 for % condition"); + annotation_error (info, + as_invalid_increment, + incr, elocus, + "increment is not constant 1 or " + "-1 for % condition"); fail = true; } } @@ -1228,9 +1356,10 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, { if (!integer_onep (i) && !integer_minus_onep (i)) { - error_at (elocus, - "increment is not constant 1 or -1 for" - " % condition"); + annotation_error (info, as_invalid_increment, + incr, elocus, + "increment is not constant 1 or -1 for" + " % condition"); fail = true; } } @@ -1242,7 +1371,8 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, } if (!incr_ok) { - error_at (elocus, "invalid increment expression"); + annotation_error (info, as_invalid_increment, incr, + elocus, "invalid increment expression"); fail = true; } } @@ -1270,6 +1400,20 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv, } } +/* External entry point to c_finish_omp_for_internal, called from the + parsers. See above for description of the arguments. */ + +tree +c_finish_omp_for (location_t locus, enum tree_code code, tree declv, + tree orig_declv, tree initv, tree condv, tree incrv, + tree body, tree pre_body, bool final_p) +{ + return c_finish_omp_for_internal (locus, code, declv, + orig_declv, initv, condv, incrv, + body, pre_body, final_p, NULL); +} + + /* Type for passing data in between c_omp_check_loop_iv and c_omp_check_loop_iv_r. */ @@ -3000,6 +3144,543 @@ c_omp_map_clause_name (tree clause, bool oacc) return omp_clause_code_name[OMP_CLAUSE_CODE (clause)]; } +/* The following functions implement automatic recognition and annotation of + for loops in OpenACC kernels regions. Inside a kernels region, a nest of + for loops that does not contain any annotated OpenACC loops, nor break + or goto statements or assignments to the variables controlling loop + termination, is converted to an OMP_FOR node with an "acc loop auto" + annotation on each loop. This feature is controlled by + flag_openacc_kernels_annotate_loops. */ + +/* Check whether DECL is the declaration of a local variable (or function + parameter) of integral type that does not have its address taken. */ + +static bool +is_local_var (tree decl) +{ + return ((TREE_CODE (decl) == VAR_DECL || TREE_CODE (decl) == PARM_DECL) + && DECL_CONTEXT (decl) != NULL + && TREE_CODE (DECL_CONTEXT (decl)) == FUNCTION_DECL + && INTEGRAL_TYPE_P (TREE_TYPE (decl)) + && !TREE_ADDRESSABLE (decl)); +} + +/* The initializer for a FOR_STMT is sometimes wrapped in various other + language-specific tree structures. We need a hook to unwrap them. + This function takes a tree argument and should return either a + MODIFY_EXPR, VAR_DECL, or NULL_TREE. */ + +static tree (*lang_specific_unwrap_initializer) (tree); + +/* Try to annotate the given NODE, which must be a FOR_STMT, with a + "#pragma acc loop auto" annotation. In practice, this means + building an OMP_FOR node for it. PREV_STMT is the statement + immediately before the loop, which may be used as the loop's + initialization statement. Annotating the loop may fail, in which + case INFO is used to record the cause of the failure and the + original loop remains unchanged. This function returns the + transformed loop if the transformation succeeded, the original node + otherwise. */ + +static tree +annotate_for_loop (tree node, tree_stmt_iterator *prev_tsi, + struct annotation_info *info) +{ + gcc_checking_assert (TREE_CODE (node) == FOR_STMT); + + location_t loc = EXPR_LOCATION (node); + tree cond = FOR_COND (node); + gcc_assert (cond); + tree decl = TREE_OPERAND (cond, 0); + gcc_assert (decl && TREE_CODE (decl) == VAR_DECL); + tree init = FOR_INIT_STMT (node); + tree prev_stmt = NULL_TREE; + bool unlink_prev = false; + bool fix_decl = false; + + + /* Both the C and C++ front ends normally put the initializer in the + statement list just before the FOR_STMT instead of in FOR_INIT_STMT. + If FOR_INIT_STMT happens to exist but isn't a MODIFY_EXPR, bail out + because the code below won't handle it. */ + if (init != NULL_TREE && TREE_CODE (init) != MODIFY_EXPR) + { + do_not_annotate_loop (info, as_invalid_initializer, NULL_TREE); + return node; + } + + /* Examine the statement before the loop to see if it is a + valid initializer. It must be either a MODIFY_EXPR or VAR_DECL, + possibly wrapped in language-specific structure. */ + if (init == NULL_TREE && prev_tsi != NULL) + { + prev_stmt = tsi_stmt (*prev_tsi); + + /* Call the language-specific hook to unwrap prev_stmt. */ + if (prev_stmt) + prev_stmt = (*lang_specific_unwrap_initializer) (prev_stmt); + + /* See if we have a valid MODIFY_EXPR. */ + if (prev_stmt + && TREE_CODE (prev_stmt) == MODIFY_EXPR + && TREE_OPERAND (prev_stmt, 0) == decl + && !TREE_SIDE_EFFECTS (TREE_OPERAND (prev_stmt, 1))) + { + init = prev_stmt; + unlink_prev = true; + } + else if (prev_stmt == decl + && !TREE_SIDE_EFFECTS (DECL_INITIAL (decl))) + { + /* If the preceding statement is the declaration of the loop + variable with its initialization, build an assignment + expression for the loop's initializer. */ + init = build2 (MODIFY_EXPR, TREE_TYPE (decl), decl, + DECL_INITIAL (decl)); + /* We need to remove the initializer from the decl if we + end up using the init we just built instead. */ + fix_decl = true; + } + } + + if (init == NULL_TREE) + /* There is nothing we can do to find the correct init statement for + this loop, but c_finish_omp_for insists on having one and would fail + otherwise. In that case, we would just return node. Do that + directly, here. */ + { + do_not_annotate_loop (info, as_missing_initializer, NULL_TREE); + return node; + } + + tree incr = FOR_EXPR (node); + + /* The C++ frontend can wrap the increment two levels deep inside a + cleanup expression, but c_finish_omp_for does not care about that. */ + if (incr != NULL_TREE && TREE_CODE (incr) == CLEANUP_POINT_EXPR) + incr = TREE_OPERAND (TREE_OPERAND (incr, 0), 0); + tree body = FOR_BODY (node); + + tree declv = make_tree_vec (1); + tree initv = make_tree_vec (1); + tree condv = make_tree_vec (1); + tree incrv = make_tree_vec (1); + TREE_VEC_ELT (declv, 0) = decl; + TREE_VEC_ELT (initv, 0) = init; + TREE_VEC_ELT (condv, 0) = cond; + TREE_VEC_ELT (incrv, 0) = incr; + + /* Do the actual transformation. This can still fail because + c_finish_omp_for has some stricter checks than we have performed up to + this point. */ + tree omp_for = c_finish_omp_for_internal (loc, OACC_LOOP, declv, NULL_TREE, + initv, condv, incrv, body, + NULL_TREE, false, info); + if (omp_for != NULL_TREE) + { + if (unlink_prev) + /* We don't need the previous statement that we consumed as an + initializer in the new OMP_FOR any more. */ + tsi_delink (prev_tsi); + + if (fix_decl) + /* We no longer need the initializer expression on the decl of + the loop variable and don't want to duplicate it. The + kernels conversion pass would interpret it as a stray + assignment in a gang-single region. */ + DECL_INITIAL (prev_stmt) = NULL_TREE; + + /* Add an auto clause, then return the new loop. */ + tree auto_clause = build_omp_clause (loc, OMP_CLAUSE_AUTO); + OMP_CLAUSE_CHAIN (auto_clause) = OMP_FOR_CLAUSES (omp_for); + OMP_FOR_CLAUSES (omp_for) = auto_clause; + return omp_for; + } + + return node; +} + +/* Forward declaration. */ +static tree annotate_loops_in_kernels_regions (tree *, int *, void *); + +/* Given a FOR_STMT NODE that is a candidate for parallelization, check its + body for validity, then try to annotate it with + "#pragma oacc loop auto", possibly modifying the current node in place. + The INFO argument contains the traversal state at the point the loop + appears. */ + +static void +check_and_annotate_for_loop (tree *nodeptr, tree_stmt_iterator *prev_tsi, + struct annotation_info *info) +{ + tree node = *nodeptr; + gcc_assert (TREE_CODE (node) == FOR_STMT); + + /* This structure describes the current loop statement. */ + struct annotation_info loop_info + = { node, NULL_TREE, false, as_in_kernels_loop, NULL_TREE, info }; + tree cond = FOR_COND (node); + + /* If we are in the body of an explicitly-annotated loop, do not add + annotations to this loop or any other nested loops. */ + if (info->state == as_explicit_annotation) + do_not_annotate_loop (&loop_info, as_explicit_annotation, info->reason); + + /* We need to find the controlling variable for the loop in order + to detect whether it is modified in the body of the loop. + That is why we are doing some checks on the loop condition + that duplicate what c_finish_omp_for is doing. */ + + /* The loop condition must be a comparison. */ + else if (cond == NULL_TREE) + do_not_annotate_loop (&loop_info, as_missing_predicate, NULL_TREE); + else if (TREE_CODE_CLASS (TREE_CODE (cond)) != tcc_comparison) + do_not_annotate_loop (&loop_info, as_invalid_predicate, cond); + else + { + /* The condition's LHS must be a local variable that does not + have its address taken. Its RHS must also be such a local + variable or a constant. */ + tree induction_var = TREE_OPERAND (cond, 0); + tree limit_var = TREE_OPERAND (cond, 1); + if (!is_local_var (induction_var) + || (!is_local_var (limit_var) + && (TREE_CODE_CLASS (TREE_CODE (limit_var)) + != tcc_constant))) + do_not_annotate_loop (&loop_info, as_invalid_predicate, cond); + else + { + /* These variables must not be assigned to in the loop. */ + loop_info.vars = tree_cons (NULL_TREE, induction_var, + loop_info.vars); + if (TREE_CODE_CLASS (TREE_CODE (limit_var)) != tcc_constant) + loop_info.vars = tree_cons (NULL_TREE, limit_var, loop_info.vars); + } + } + + /* Walk the body. This will process any nested loops, so we have to do it + even if we have already rejected this loop as a candidate for + annotation. */ + walk_tree (&FOR_BODY (node), annotate_loops_in_kernels_regions, + (void *) &loop_info, NULL); + + if (loop_info.state == as_in_kernels_loop) + { + /* If the traversal of the loop and all nested loops didn't hit + any problems, attempt the actual transformation. If it + succeeds, replace this node with the annotated loop. */ + tree result = annotate_for_loop (node, prev_tsi, &loop_info); + if (result != node) + { + /* Success! */ + *nodeptr = result; + return; + } + } + + /* If we got here, we have a FOR_STMT we could not convert to an + OMP loop. */ + + if (loop_info.state == as_invalid_return) + /* This is diagnosed elsewhere as a hard error, so no warning is + needed here. */ + return; + + /* Issue warnings about other problems. */ + auto_diagnostic_group d; + if (warning_at (EXPR_LOCATION (node), + OPT_Wopenacc_kernels_annotate_loops, + "loop cannot be annotated for OpenACC parallelization")) + { + location_t locus; + if (loop_info.reason && EXPR_HAS_LOCATION (loop_info.reason)) + locus = EXPR_LOCATION (loop_info.reason); + else + locus = EXPR_LOCATION (node); + switch (loop_info.state) + { + case as_invalid_variable_type: + inform (locus, "invalid type for iteration variable %qE", + loop_info.reason); + break; + case as_missing_initializer: + inform (locus, "missing iteration variable initializer"); + break; + case as_invalid_initializer: + inform (locus, "unrecognized initializer"); + break; + case as_missing_predicate: + inform (locus, "missing controlling predicate"); + break; + case as_invalid_predicate: + inform (locus, "invalid controlling predicate"); + break; + case as_missing_increment: + inform (locus, "missing increment expression"); + break; + case as_invalid_increment: + inform (locus, "invalid increment expression"); + break; + case as_explicit_annotation: + inform (locus, "explicit OpenACC annotation in loop nest"); + break; + case as_invalid_control_flow: + inform (locus, "loop contains unstructured control flow"); + break; + case as_invalid_break: + inform (locus, "loop contains % statement"); + break; + case as_invalid_call: + inform (locus, "loop contains call to non-oacc function"); + break; + case as_invalid_modification: + inform (locus, "invalid modification of controlling variable"); + break; + default: + gcc_unreachable (); + } + } +} + +/* Traversal function for walk_tree. Visit the tree, finding OpenACC + kernels regions. DATA is NULL if we are outside of a kernels region, + otherwise it is a pointer to the enclosing kernels region's + annotation_info struct. If the traversal encounters a for loop inside a + kernels region that is a candidate for parallelization, annotate it + with OpenACC loop directives. */ + +static tree +annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees, + void *data) +{ + tree node = *nodeptr; + struct annotation_info *info = (struct annotation_info *) data; + gcc_assert (info); + + switch (TREE_CODE (node)) + { + case OACC_KERNELS: + /* Recursively process the body of the kernels region in a new info + scope. */ + if (info->state == as_outer) + { + struct annotation_info nested_info + = { NULL_TREE, NULL_TREE, true, + as_in_kernels_region, NULL_TREE, info }; + walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, + (void *) &nested_info, NULL); + *walk_subtrees = 0; + } + break; + + case OACC_LOOP: + /* Do not try to add automatic OpenACC annotations inside manually + annotated loops. Presumably, the user avoided doing it on + purpose; for example, all available levels of parallelism may + have been used up. */ + { + struct annotation_info nested_info + = { NULL_TREE, NULL_TREE, false, as_explicit_annotation, + node, info }; + if (info->state >= as_in_kernels_region) + do_not_annotate_loop_nest (info, as_explicit_annotation, + node); + walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, + (void *) &nested_info, NULL); + *walk_subtrees = 0; + } + break; + + case FOR_STMT: + /* Try to annotate the loop if we are in a kernels region. + This will do a recursive traversal of the loop body in a new + info scope. */ + if (info->state >= as_in_kernels_region) + { + check_and_annotate_for_loop (nodeptr, NULL, info); + *walk_subtrees = 0; + } + break; + + case LABEL_EXPR: + /* Possibly unstructured control flow. Unless we perform further + analyses, we must assume that such control flow may enter the + current loop. In this case, we must not parallelize the loop. */ + if (info->state >= as_in_kernels_loop + && TREE_USED (LABEL_EXPR_LABEL (node))) + do_not_annotate_loop_nest (info, as_invalid_control_flow, node); + break; + + case GOTO_EXPR: + /* Possibly unstructured control flow. Unless we perform further + analyses, we must assume that such control flow may leave the + current loop. In this case, we must not parallelize the loop. */ + if (info->state >= as_in_kernels_loop) + do_not_annotate_loop_nest (info, as_invalid_control_flow, node); + break; + + case BREAK_STMT: + /* A break statement. Whether or not this is valid depends on the + enclosing context. */ + if (info->state >= as_in_kernels_loop && !info->break_ok) + do_not_annotate_loop (info, as_invalid_break, node); + break; + + case RETURN_EXPR: + /* A return leaves the entire loop nest. */ + if (info->state >= as_in_kernels_loop) + do_not_annotate_loop_nest (info, as_invalid_return, node); + break; + + case CALL_EXPR: + /* Direct function calls to functions marked as OpenACC routines are + allowed. Reject indirect calls or calls to non-routines. */ + if (info->state >= as_in_kernels_loop) + { + tree fn = CALL_EXPR_FN (node), fn_decl = NULL_TREE; + if (fn != NULL_TREE && TREE_CODE (fn) == FUNCTION_DECL) + fn_decl = fn; + else if (fn != NULL_TREE && TREE_CODE (fn) == ADDR_EXPR) + { + tree fn_op = TREE_OPERAND (fn, 0); + if (fn_op != NULL_TREE && TREE_CODE (fn_op) == FUNCTION_DECL) + fn_decl = fn_op; + } + if (fn_decl == NULL_TREE) + do_not_annotate_loop_nest (info, as_invalid_call, node); + else if (!lookup_attribute ("oacc function", + DECL_ATTRIBUTES (fn_decl))) + do_not_annotate_loop_nest (info, as_invalid_call, node); + } + break; + + case MODIFY_EXPR: + /* See if this assignment's LHS is one of the variables that must + not be modified in the loop body because they control termination + of the loop (or an enclosing loop in the nest). */ + if (info->state >= as_in_kernels_loop) + { + tree lhs = TREE_OPERAND (node, 0); + if (!is_local_var (lhs)) + /* Early exit: This cannot be a variable we care about. */ + break; + /* Walk up the loop stack. Invalidate the ones controlled by this + variable. There may be several, if this variable is the common + iteration limit for several nested loops. */ + for (struct annotation_info *outer_loop = info; outer_loop != NULL; + outer_loop = outer_loop->next) + for (tree t = outer_loop->vars; t != NULL_TREE; t = TREE_CHAIN (t)) + if (TREE_VALUE (t) == lhs) + { + do_not_annotate_loop (outer_loop, + as_invalid_modification, + node); + break; + } + } + break; + + case SWITCH_STMT: + /* Needs special handling to allow break in the body. */ + if (info->state >= as_in_kernels_loop) + { + bool save_break_ok = info->break_ok; + + walk_tree (&SWITCH_STMT_COND (node), + annotate_loops_in_kernels_regions, + (void *) info, NULL); + info->break_ok = true; + walk_tree (&SWITCH_STMT_BODY (node), + annotate_loops_in_kernels_regions, + (void *) info, NULL); + info->break_ok = save_break_ok; + *walk_subtrees = 0; + } + break; + + case WHILE_STMT: + /* Needs special handling to allow break in the body. */ + if (info->state >= as_in_kernels_loop) + { + bool save_break_ok = info->break_ok; + + walk_tree (&WHILE_COND (node), annotate_loops_in_kernels_regions, + (void *) info, NULL); + info->break_ok = true; + walk_tree (&WHILE_BODY (node), annotate_loops_in_kernels_regions, + (void *) info, NULL); + info->break_ok = save_break_ok; + *walk_subtrees = 0; + } + break; + + case DO_STMT: + /* Needs special handling to allow break in the body. */ + if (info->state >= as_in_kernels_loop) + { + bool save_break_ok = info->break_ok; + + walk_tree (&DO_COND (node), annotate_loops_in_kernels_regions, + (void *) info, NULL); + info->break_ok = true; + walk_tree (&DO_BODY (node), annotate_loops_in_kernels_regions, + (void *) info, NULL); + info->break_ok = save_break_ok; + *walk_subtrees = 0; + } + break; + + case STATEMENT_LIST: + /* We iterate over these explicitly so that we can track the previous + statement in the chain. It may be the initializer for a following + FOR_STMT node. */ + if (info->state >= as_in_kernels_region) + { + tree_stmt_iterator i = tsi_start (node); + tree_stmt_iterator prev, *prev_tsi = NULL; + while (!tsi_end_p (i)) + { + tree *stmtptr = tsi_stmt_ptr (i); + if (TREE_CODE (*stmtptr) == FOR_STMT) + { + check_and_annotate_for_loop (stmtptr, prev_tsi, info); + *walk_subtrees = 0; + } + else + walk_tree (stmtptr, annotate_loops_in_kernels_regions, + (void *) info, NULL); + prev = i; + prev_tsi = &prev; + tsi_next (&i); + } + *walk_subtrees = 0; + } + break; + + default: + break; + } + + return NULL_TREE; +} + +/* Find for loops in OpenACC kernels regions that do not have OpenACC + annotations but look like they might benefit from automatic + parallelization. Convert them from FOR_STMT to OMP_FOR nodes and + add the equivalent of "#pragma acc loop auto" annotations for them. + Assumes flag_openacc_kernels_annotate_loops is set. */ + +void +c_oacc_annotate_loops_in_kernels_regions (tree decl, + tree (*unwrap_fn) (tree)) +{ + struct annotation_info info + = { NULL_TREE, NULL_TREE, true, as_outer, NULL_TREE, NULL }; + lang_specific_unwrap_initializer = unwrap_fn; + walk_tree (&DECL_SAVED_TREE (decl), annotate_loops_in_kernels_regions, + (void *) &info, NULL); +} + /* Used to merge map clause information in c_omp_adjust_map_clauses. */ struct map_clause { diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt index 06457ac739e4..a0f43d6d325f 100644 --- a/gcc/c-family/c.opt +++ b/gcc/c-family/c.opt @@ -1074,6 +1074,10 @@ Wopenacc-parallelism C C++ Var(warn_openacc_parallelism) Warning Warn about potentially suboptimal choices related to OpenACC parallelism. +Wopenacc-kernels-annotate-loops +C ObjC C++ ObjC++ Warning Var(warn_openacc_kernels_annotate_loops) Init(0) +Warn about loops in OpenACC kernels regions that cannot be parallelized. + Wopenmp-simd C C++ Var(warn_openmp_simd) Warning LangEnabledBy(C C++,Wall) Warn if a simd directive is overridden by the vectorizer cost model. @@ -1910,6 +1914,10 @@ fopenacc-dim= C ObjC C++ ObjC++ LTO Joined Var(flag_openacc_dims) Specify default OpenACC compute dimensions. +fopenacc-kernels-annotate-loops +C ObjC C++ ObjC++ LTO Optimization Var(flag_openacc_kernels_annotate_loops) Init(1) +Automatically parallelize unannotated loops in OpenACC kernels regions. + fopenmp C ObjC C++ ObjC++ LTO Var(flag_openmp) Enable OpenMP (implies -frecursive in Fortran). diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c index 186fa1692c16..467b3425b9a4 100644 --- a/gcc/c/c-decl.c +++ b/gcc/c/c-decl.c @@ -10230,6 +10230,29 @@ temp_pop_parm_decls (void) pop_scope (); } +/* Function passed to c_oacc_annotate_loop_in_kernels_regions to do + language-specific unwrapping of an initializer expression. */ +static tree +c_unwrap_for_init (tree x) +{ + if (!x) + return NULL_TREE; + + while (true) + switch (TREE_CODE (x)) + { + case MODIFY_EXPR: + case VAR_DECL: + return x; + + case DECL_EXPR: + x = TREE_OPERAND (x, 0); + break; + + default: + return NULL_TREE; + } +} /* Finish up a function declaration and compile that function all the way to assembler language output. Then free the storage @@ -10332,6 +10355,11 @@ finish_function (location_t end_loc) if (warn_unused_parameter) do_warn_unused_parameter (fndecl); + /* If requested, automatically annotate suitable loops in OpenACC kernels + regions with OpenACC loop annotations to allow auto-parallelization. */ + if (flag_openacc && flag_openacc_kernels_annotate_loops) + c_oacc_annotate_loops_in_kernels_regions (fndecl, c_unwrap_for_init); + /* Store the end of the function, so that we get good line number info for the epilogue. */ cfun->function_end_locus = end_loc; diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 7c2a134e4061..17f14d1f6742 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -17528,6 +17528,45 @@ emit_coro_helper (tree helper) expand_or_defer_fn (helper); } + +/* Function passed to c_oacc_annotate_loop_in_kernels_regions to do + language-specific unwrapping of an initializer expression. */ +static tree +cp_unwrap_for_init (tree x) +{ + if (!x) + return NULL_TREE; + + while (true) + switch (TREE_CODE (x)) + { + case MODIFY_EXPR: + case VAR_DECL: + return x; + + case CLEANUP_POINT_EXPR: + x = TREE_OPERAND (x, 0); + break; + + case EXPR_STMT: + x = TREE_OPERAND (x, 0); + break; + + case DECL_EXPR: + x = TREE_OPERAND (x, 0); + break; + + case CONVERT_EXPR: + if (TREE_TYPE (x) != void_type_node) + return NULL_TREE; + x = TREE_OPERAND (x, 0); + break; + + default: + return NULL_TREE; + } +} + /* Finish up a function declaration and compile that function all the way to assembler language output. The free the storage for the function definition. INLINE_P is TRUE if we just @@ -17832,6 +17871,11 @@ finish_function (bool inline_p) && !DECL_CLONED_FUNCTION_P (fndecl)) do_warn_unused_parameter (fndecl); + /* If requested, automatically annotate suitable loops in OpenACC kernels + regions with OpenACC loop annotations to allow auto-parallelization. */ + if (flag_openacc && flag_openacc_kernels_annotate_loops) + c_oacc_annotate_loops_in_kernels_regions (fndecl, cp_unwrap_for_init); + /* Genericize before inlining. */ if (!processing_template_decl && !DECL_IMMEDIATE_FUNCTION_P (fndecl) diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index 9fb74d349203..e0f09610408c 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -371,6 +371,7 @@ Objective-C and Objective-C++ Dialects}. -Wnull-dereference -Wno-odr @gol -Wopenacc-parallelism @gol -Wopenmp-simd @gol +-Wopenacc-kernels-annotate-loops -Wopenmp-simd @gol -Wno-overflow -Woverlength-strings -Wno-override-init-side-effects @gol -Wpacked -Wno-packed-bitfield-compat -Wpacked-not-aligned -Wpadded @gol -Wparentheses -Wno-pedantic-ms-format @gol @@ -533,7 +534,8 @@ Objective-C and Objective-C++ Dialects}. -fmerge-constants -fmodulo-sched -fmodulo-sched-allow-regmoves @gol -fmove-loop-invariants -fmove-loop-stores -fno-branch-count-reg @gol -fno-defer-pop -fno-fp-int-builtin-inexact -fno-function-cse @gol --fno-guess-branch-probability -fno-inline -fno-math-errno -fno-peephole @gol +-fno-guess-branch-probability -fno-inline -fno-math-errno @gol +-fno-openacc-kernels-annotate-loops -fno-peephole @gol -fno-peephole2 -fno-printf-return-value -fno-sched-interblock @gol -fno-sched-spec -fno-signed-zeros @gol -fno-toplevel-reorder -fno-trapping-math -fno-zero-initialized-in-bss @gol @@ -8957,6 +8959,13 @@ Enabled by default. @cindex OpenACC accelerator programming Warn about potentially suboptimal choices related to OpenACC parallelism. +@item -Wopenacc-kernels-annotate-loops +@opindex Wopenacc-kernels-annotate-loops +@opindex Wno-Wopenacc-kernels-annotate-loops +Warn about @code{for} (C/C++) or @code{DO} (Fortran) loops in OpenACC +kernels regions that cannot be automatically annotated for +parallelization with @option{-fopenacc-kernels-annotate-loops}. + @item -Wopenmp-simd @opindex Wopenmp-simd @opindex Wno-openmp-simd @@ -14835,6 +14844,27 @@ SIMD iterations. @end table +@item -fno-openacc-kernels-annotate-loops +@opindex fno-openacc-kernels-annotate-loops +@opindex fopenacc-kernels-annotate-loops +@cindex kernels regions, OpenACC +Disable automatic parallelization of unannotated loops in OpenACC +kernels regions. The default is to attempt to add implicit +@code{acc loop auto} annotations to loops in kernels regions if +@option{-fopenacc} is enabled. + +Note that you can use @option{-Wopenacc-kernels-annotate-loops} to +diagnose @code{for} loops that cannot be automatically annotated +(@pxref{Warning Options}). Reasons why automatic loop annotations +cannot be applied include premature exits, calls to functions without +an @code{openacc routine} annotation, or unstructured control flow in +the loop body. In C and C++, the loop variable initialization, end +test, and increment expressions must additionally conform to +restrictions similar to those for explicitly-annotated loops, and the +loop variable must not be otherwise modified in the body of the loop. +An explicit @code{acc loop} annotation disables automatic annotations +on any nested or containing loops. + @end table @node Instrumentation Options 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 1d12658790d1..e391184f403d 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -2,6 +2,7 @@ OpenACC kernels. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fno-openacc-kernels-annotate-loops" } { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index bdf7b4a06410..779e2b0a24db 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -1,7 +1,8 @@ /* Check offloaded function's attributes and classification for OpenACC - kernels. */ + 'kernels' (parloops version). */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fno-openacc-kernels-annotate-loops" } { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c index 030425475495..c37152c74041 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-dom3" } */ #include diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-counter-vars-function-scope.c b/gcc/testsuite/c-c++-common/goacc/kernels-counter-vars-function-scope.c index c475333f1aef..b1f43029af7c 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-counter-vars-function-scope.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-counter-vars-function-scope.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c index 8f7f415b58d8..e87aab3295c7 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c index c11d36fb4373..2323857fb4ad 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c index acef6a1a1793..adca30bf2cd7 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-2.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c index 75e2bb78cea4..5f16085ff386 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-3.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-1.c new file mode 100644 index 000000000000..c7b5ac882195 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-1.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that all loops in the nest are annotated. */ + +void f (float a[16][16], float b[16][16], float c[16][16]) +{ + int i, j, k; + +#pragma acc kernels copyin(a[0:16][0:16], b[0:16][0:16]) copyout(c[0:16][0:16]) + { + for (i = 0; i < 16; i++) { + for (j = 0; j < 16; j++) { + float t = 0; + for (k = 0; k < 16; k++) + t += a[i][k] * b[k][j]; + c[i][j] = t; + } + } + } + +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 3 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-10.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-10.c new file mode 100644 index 000000000000..58b41d20e232 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-10.c @@ -0,0 +1,32 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a random goto in the body can't be annotated. */ + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + { + if (a[i] < 0) + { + t = 0; + goto bad; + } + t += a[i] * b[i]; + } + bad: + ; + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-11.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-11.c new file mode 100644 index 000000000000..e9d2ef48611a --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-11.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a random label in the body triggers a warning. */ + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i = n - 1; + +#pragma acc kernels + { + goto spaghetti; + for (i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + { + spaghetti: + t += a[i] * b[i]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-12.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-12.c new file mode 100644 index 000000000000..ba408bc3634d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-12.c @@ -0,0 +1,28 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that in a situation with nested loops, a problem that prevents + annotation of the inner loop only still allows the outer loop to be + annotated. */ + +float f (float *a, float *b, int n) +{ + float t = 0; + +#pragma acc kernels + { + for (int i = 0; i < n; i++) + for (int j = 0; j <= i; j++) /* { dg-warning "loop cannot be annotated" } */ + { + if (a[i] < 0 || b[j] < 0) + j = i; + else + t += a[i] * b[j]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-13.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-13.c new file mode 100644 index 000000000000..64433e816ed4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-13.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that in a situation with nested loops, a problem that prevents + annotation of the outer loop only still allows the inner loop to be + annotated. */ + +float f (float *a, float *b, int n) +{ + float t = 0; + +#pragma acc kernels + { + for (int i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + { + if (a[i] < 0) + n = i; + for (int j = 0; j <= i; j++) + t += a[i] * b[j]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-14.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-14.c new file mode 100644 index 000000000000..379e6baf97c3 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-14.c @@ -0,0 +1,22 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that an explicit annotation on an outer loop suppresses annotation + of inner loops, and produces a diagnostic. */ + +void f (float *a, float *b) +{ + float t = 0; + +#pragma acc kernels + { +#pragma acc loop seq + for (int l = 0; l < 20; l++) + for (int m = 0; m < 20; m++) /* { dg-warning "loop cannot be annotated" } */ + b[m] = a[m]; + } +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-15.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-15.c new file mode 100644 index 000000000000..9a2a7cabde5d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-15.c @@ -0,0 +1,22 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that an explicit annotation on an inner loop suppresses annotation + of outer loops, and produces a diagnostic. */ + +void f (float *a, float *b) +{ + float t = 0; + +#pragma acc kernels + { + for (int l = 0; l < 20; l++) /* { dg-warning "loop cannot be annotated" } */ +#pragma acc loop seq + for (int m = 0; m < 20; m++) + b[m] = a[m]; + } +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-16.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-16.c new file mode 100644 index 000000000000..075f897fad4a --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-16.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a modification of the loop variable in the + body cannot be annotated. */ + +float f (float *a, float *b, int n) +{ + float t = 0; + +#pragma acc kernels + { + for (int i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + { + if (a[i] < 0 || b[i] < 0) + i = n; + else + t += a[i] * b[i]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-17.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-17.c new file mode 100644 index 000000000000..507678965b4d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-17.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a modification of the loop iteration count + variable in the body cannot be annotated. */ + +float f (float *a, float *b, int n) +{ + float t = 0; + +#pragma acc kernels + { + for (int i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + { + if (a[i] < 0 || b[i] < 0) + n = i; + else + t += a[i] * b[i]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-2.c new file mode 100644 index 000000000000..9e0a946828ff --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-2.c @@ -0,0 +1,21 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a variable bound can be annotated. */ + +float f (float *a, float *b, int n) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) + t += a[i] * b[i]; + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-3.c new file mode 100644 index 000000000000..f60070e27961 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-3.c @@ -0,0 +1,24 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a conditional in the body can be annotated. */ + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) + if (a[i] > 0 && b[i] > 0) + t += a[i] * b[i]; + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-4.c new file mode 100644 index 000000000000..949871cc42ec --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-4.c @@ -0,0 +1,34 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a switch and break in the body can be annotated. */ + +#define n 16 + +float f (float *a, float *b, int state) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) + switch (state) + { + case 0: + default: + t += a[i] * b[i]; + break; + + case 1: + if (a[i] > 0 && b[i] > 0) + t += a[i] * b[i]; + break; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-5.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-5.c new file mode 100644 index 000000000000..03dfe8fbcd40 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-5.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a continue statement in the body can be annotated. */ + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) + { + if (a[i] < 0 || b[i] < 0) + continue; + t += a[i] * b[i]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-6.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-6.c new file mode 100644 index 000000000000..ede6b3c8cd67 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-6.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a break statement in the body cannot be annotated. */ + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + { + if (a[i] < 0 || b[i] < 0) + break; + t += a[i] * b[i]; + } + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-7.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-7.c new file mode 100644 index 000000000000..20ee29989665 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-7.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with a random function call in the body cannot be + annotated. */ + +extern float g (float); + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) /* { dg-warning "loop cannot be annotated" } */ + t += g (a[i] * b[i]); + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-8.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-8.c new file mode 100644 index 000000000000..796f048d67ca --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-8.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a loop with an openacc function call in the body can be + annotated. */ + +#pragma acc routine worker +extern float g (float); + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) + t += g (a[i] * b[i]); + } + return t; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-9.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-9.c new file mode 100644 index 000000000000..048f1b09a84d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-9.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-do compile } */ + +/* Test that a kernels loop with a return in the body triggers a hard + error. */ + +#define n 16 + +float f (float *a, float *b) +{ + float t = 0; + int i; + +#pragma acc kernels + { + for (i = 0; i < n; i++) + { + if (a[i] < 0 || b[i] < 0) + return 0.0; /* { dg-error "invalid branch" } */ + t += a[i] * b[i]; + } + } + return t; +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-2.c index 71800217991a..9a97de6f6e13 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-2.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit-2.c index 0c9f83312408..31e8378e3d74 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit-2.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit.c index 0bd21b68d317..ad591551b979 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-enter-exit.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-update.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-update.c index dd5a84146a8e..4acffef41ba1 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-update.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data-update.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data.c index a658182de904..327aa0570c9c 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-data.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-data.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c index 73b469d70610..26c65fe742aa 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c @@ -1,5 +1,6 @@ /* { dg-additional-options "-O2" } */ /* { dg-additional-options "-g" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-mod-not-zero.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-mod-not-zero.c index 55926230d578..8955cf29224b 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-mod-not-zero.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-mod-not-zero.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c index e86be1b1cdc0..d88a61dbab51 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-n.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c index 2b0e186ae297..5943d56a5bbe 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-nest.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop.c index 9619d53b43d7..ad525cdbe141 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-one-counter-var.c b/gcc/testsuite/c-c++-common/goacc/kernels-one-counter-var.c index 69539b24a78d..f799baffd8df 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-one-counter-var.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-one-counter-var.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c b/gcc/testsuite/c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c index 81b0fee5a44c..b8093b54dec8 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-parallel-loop-data-enter-exit.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-reduction.c b/gcc/testsuite/c-c++-common/goacc/kernels-reduction.c index 5921b88920fd..105cbcf3ba2e 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-reduction.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-reduction.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fno-openacc-kernels-annotate-loops" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */