From patchwork Wed Dec 15 15:54:10 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48945 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 6C2133858430 for ; Wed, 15 Dec 2021 15:57:27 +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 9D6E2385801A; Wed, 15 Dec 2021 15:55:07 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 9D6E2385801A 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: Kn6E/YZmsiVfD1m8lFen7uHEt6nQP/59sL00Owq9DQk66RvGfZziER1O3v6DPuqAzyeiX2D8ub OfEdkTDaznQ7uCVySdHUp/Aia86hyY87eLOye1m6DzLqyWlBLd0ls3dqsZ8UujZ3Utou9jko2O r6dP3ujK5FM3Tz9qHZ3wnnhC3E5JyYdYfa30BL6iPZJS7lriz8zIwTFLVlOB8yzUg7JbX3UDOg qOI7ebh44I8uiBkZTKcptiG3gGy1CFrFRLZhPrTIhYoVsXQPSJqam0uy2YrNLaAf9KY+3+Q+T8 b99DHWatyvDqG4gtBGxbHKmG X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69584566" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:07 -0800 IronPort-SDR: BGzLLOwbJAxKbvdTntHJSj1qjSYsxJ7WHy+jaxwyCq5yjsYwpClM+9QGkSMjheYWo+/JWPniNJ VVC2aboaQFEknfG/jUswsZFbhiJyyhBiy8BmhkiI+CaQDMvh+fwBZxpXzN/2K8E1pJWiHX/xAD IeJtmplW4CF1wnJiHjoIHeccnRgEfGR158LLfJMHKEitumW81AXVnMJ5lKzlM6UlIXN1ZtlAZs NyOp/bdD1D8Pj49tW4965QqhjLME3pfwa1qvQ4ixNGqTGxMI6AZc1dJcsfeINQEoCZh8u/ti2r YJw= From: Frederik Harwath To: Subject: [PATCH 03/40] Kernels loops annotation: Fortran. Date: Wed, 15 Dec 2021 16:54:10 +0100 Message-ID: <20211215155447.19379-4-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=-12.5 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: tobias@codesourcery.com, Sandra Loosemore , =?utf-8?q?Gerg=C3=B6_Barany?= , thomas@codesourcery.com, fortran@gcc.gnu.org Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" From: Sandra Loosemore This patch implements the Fortran support for adding "#pragma acc loop auto" annotations to loops in OpenACC kernels regions. It implements the same -fopenacc-kernels-annotate-loops and -Wopenacc-kernels-annotate-loops options that were previously added (and documented) for the C/C++ front ends. Co-Authored-By: Gergö Barany gcc/fortran/ * gfortran.h (gfc_oacc_annotate_loops_in_kernels_regions): Declare. * lang.opt (Wopenacc-kernels-annotate-loops): New. (fopenacc-kernels-annotate-loops): New. * openmp.c: Include options.h. (enum annotation_state, enum annotation_result): New. (check_code_for_invalid_calls): New. (check_expr_for_invalid_calls): New. (check_for_invalid_calls): New. (annotate_do_loop): New. (annotate_do_loops_in_kernels): New. (compute_goto_targets): New. (gfc_oacc_annotate_loops_in_kernels_regions): New. * parse.c (gfc_parse_file): Handle -fopenacc-kernels-annotate-loops. gcc/testsuite/ * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add -fno-openacc-kernels-annotate-loops option. * gfortran.dg/goacc/classify-kernels.f95: Likewise. * gfortran.dg/goacc/common-block-3.f90: Likewise. * gfortran.dg/goacc/kernels-loop-2.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-2.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-update.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data.f95: Likewise. * gfortran.dg/goacc/kernels-loop-n.f95: Likewise. * gfortran.dg/goacc/kernels-loop.f95: Likewise. * gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95: Likewise. * gfortran.dg/goacc/kernels-loop-annotation-1.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-2.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-3.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-4.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-5.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-6.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-7.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-8.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-9.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-10.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-11.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-12.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-13.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-14.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-15.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-16.f95: New. --- gcc/fortran/gfortran.h | 1 + gcc/fortran/lang.opt | 8 + gcc/fortran/openmp.c | 364 ++++++++++++++++++ gcc/fortran/parse.c | 9 + .../goacc/classify-kernels-unparallelized.f95 | 1 + .../gfortran.dg/goacc/classify-kernels.f95 | 1 + .../gfortran.dg/goacc/common-block-3.f90 | 1 + .../gfortran.dg/goacc/kernels-loop-2.f95 | 1 + .../goacc/kernels-loop-annotation-1.f95 | 33 ++ .../goacc/kernels-loop-annotation-10.f95 | 32 ++ .../goacc/kernels-loop-annotation-11.f95 | 34 ++ .../goacc/kernels-loop-annotation-12.f95 | 39 ++ .../goacc/kernels-loop-annotation-13.f95 | 38 ++ .../goacc/kernels-loop-annotation-14.f95 | 35 ++ .../goacc/kernels-loop-annotation-15.f95 | 35 ++ .../goacc/kernels-loop-annotation-16.f95 | 34 ++ .../goacc/kernels-loop-annotation-2.f95 | 32 ++ .../goacc/kernels-loop-annotation-3.f95 | 33 ++ .../goacc/kernels-loop-annotation-4.f95 | 34 ++ .../goacc/kernels-loop-annotation-5.f95 | 35 ++ .../goacc/kernels-loop-annotation-6.f95 | 34 ++ .../goacc/kernels-loop-annotation-7.f95 | 48 +++ .../goacc/kernels-loop-annotation-8.f95 | 50 +++ .../goacc/kernels-loop-annotation-9.f95 | 34 ++ .../gfortran.dg/goacc/kernels-loop-data-2.f95 | 1 + .../goacc/kernels-loop-data-enter-exit-2.f95 | 1 + .../goacc/kernels-loop-data-enter-exit.f95 | 1 + .../goacc/kernels-loop-data-update.f95 | 1 + .../gfortran.dg/goacc/kernels-loop-data.f95 | 1 + .../gfortran.dg/goacc/kernels-loop-n.f95 | 1 + .../gfortran.dg/goacc/kernels-loop.f95 | 1 + .../kernels-parallel-loop-data-enter-exit.f95 | 1 + 32 files changed, 974 insertions(+) create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95 -- 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/fortran/gfortran.h b/gcc/fortran/gfortran.h index f7662c59a5df..50db768ce0fc 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -3545,6 +3545,7 @@ void gfc_resolve_oacc_declare (gfc_namespace *); void gfc_resolve_oacc_parallel_loop_blocks (gfc_code *, gfc_namespace *); void gfc_resolve_oacc_blocks (gfc_code *, gfc_namespace *); void gfc_resolve_oacc_routines (gfc_namespace *); +void gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *); /* expr.c */ void gfc_free_actual_arglist (gfc_actual_arglist *); diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt index 6db01c736be1..a202c04c4a25 100644 --- a/gcc/fortran/lang.opt +++ b/gcc/fortran/lang.opt @@ -289,6 +289,10 @@ Wopenacc-parallelism Fortran ; Documented in C +Wopenacc-kernels-annotate-loops +Fortran +; Documented in C + Wopenmp-simd Fortran ; Documented in C @@ -695,6 +699,10 @@ fopenacc-dim= Fortran LTO Joined Var(flag_openacc_dims) ; Documented in C +fopenacc-kernels-annotate-loops +Fortran LTO Optimization +; Documented in C + fopenmp Fortran LTO ; Documented in C diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index dcf22ac2c2f3..243b5e0a9ac6 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -29,6 +29,7 @@ along with GCC; see the file COPYING3. If not see #include "diagnostic.h" #include "gomp-constants.h" #include "target-memory.h" /* For gfc_encode_character. */ +#include "options.h" /* Match an end of OpenMP directive. End of OpenMP directive is optional whitespace, followed by '\n' or comment '!'. */ @@ -9090,3 +9091,366 @@ gfc_resolve_omp_udrs (gfc_symtree *st) for (omp_udr = st->n.omp_udr; omp_udr; omp_udr = omp_udr->next) gfc_resolve_omp_udr (omp_udr); } + + +/* The following functions implement automatic recognition and annotation of + DO loops in OpenACC kernels regions. Inside a kernels region, a nest of + DO loops that does not contain any annotated OpenACC loops, nor EXIT + or GOTO statements, gets an automatic "acc loop auto" annotation + on each loop. + This feature is controlled by flag_openacc_kernels_annotate_loops. */ + + +/* State of annotation state traversal for DO loops in kernels regions. */ +enum annotation_state { + as_outer, + as_in_kernels_region, + as_in_kernels_loop, + as_in_kernels_inner_loop +}; + +/* Return status of annotation traversal. */ +enum annotation_result { + ar_ok, + ar_invalid_loop, + ar_invalid_nest +}; + +/* Code walk function for check_for_invalid_calls. */ + +static int +check_code_for_invalid_calls (gfc_code **codep, int *walk_subtrees, + void *data ATTRIBUTE_UNUSED) +{ + gfc_code *code = *codep; + switch (code->op) + { + case EXEC_CALL: + /* Calls to openacc routines are permitted. */ + if (code->resolved_sym + && (code->resolved_sym->attr.oacc_routine_lop + != OACC_ROUTINE_LOP_NONE)) + return 0; + /* Else fall through. */ + + case EXEC_CALL_PPC: + case EXEC_ASSIGN_CALL: + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "Subroutine call at %L prevents annotation of loop nest", + &code->loc); + *walk_subtrees = 0; + return 1; + + default: + return 0; + } +} + +/* Expr walk function for check_for_invalid_calls. */ + +static int +check_expr_for_invalid_calls (gfc_expr **exprp, int *walk_subtrees, + void *data ATTRIBUTE_UNUSED) +{ + gfc_expr *expr = *exprp; + switch (expr->expr_type) + { + case EXPR_FUNCTION: + if (expr->value.function.esym + && (expr->value.function.esym->attr.oacc_routine_lop + != OACC_ROUTINE_LOP_NONE)) + return 0; + /* Else fall through. */ + + case EXPR_COMPCALL: + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "Function call at %L prevents annotation of loop nest", + &expr->where); + *walk_subtrees = 0; + return 1; + + default: + return 0; + } +} + +/* Return TRUE if the DO loop CODE contains function or procedure + calls that ought to prohibit annotation. This traversal is + separate from the main annotation tree walk because we need to walk + expressions as well as executable statements. */ + +static bool +check_for_invalid_calls (gfc_code *code) +{ + gcc_assert (code->op == EXEC_DO); + return gfc_code_walker (&code, check_code_for_invalid_calls, + check_expr_for_invalid_calls, NULL); +} + +/* Annotate DO loop CODE with OpenACC "loop auto". */ + +static void +annotate_do_loop (gfc_code *code, gfc_code *parent) +{ + + /* A DO loop's body is another phony DO node whose next pointer starts + the actual body. */ + gcc_assert (code->op == EXEC_DO); + gcc_assert (code->block->op == EXEC_DO); + + /* Build the "acc loop auto" annotation and add the loop as its + body. */ + gfc_omp_clauses *clauses = gfc_get_omp_clauses (); + clauses->par_auto = 1; + gfc_code *oacc_loop = gfc_get_code (EXEC_OACC_LOOP); + oacc_loop->block = gfc_get_code (EXEC_OACC_LOOP); + oacc_loop->block->next = code; + oacc_loop->ext.omp_clauses = clauses; + oacc_loop->loc = code->loc; + oacc_loop->block->loc = code->loc; + + /* Splice the annotation into the place of the original loop. */ + if (parent->block == code) + parent->block = oacc_loop; + else + { + gfc_code *prev = parent->block; + while (prev != code && prev->next != code) + { + prev = prev->next; + gcc_assert (prev != NULL); + } + prev->next = oacc_loop; + } + oacc_loop->next = code->next; + code->next = NULL; +} + +/* Recursively traverse CODE in block PARENT, finding OpenACC kernels + regions. GOTO_TARGETS keeps track of statement labels that are + targets of gotos in the current function, while STATE keeps track + of the current context of the traversal. If the traversal + encounters a DO loop inside a kernels region, annotate it with + OpenACC loop directives if appropriate. Return the status of the + traversal. */ + +static enum annotation_result +annotate_do_loops_in_kernels (gfc_code *code, gfc_code *parent, + hash_set *goto_targets, + annotation_state state) +{ + gfc_code *next_code = NULL; + enum annotation_result retval = ar_ok; + + for ( ; code; code = next_code) + { + bool walk_block = true; + next_code = code->next; + + if (state >= as_in_kernels_loop + && code->here && goto_targets->contains (code->here)) + /* This statement has a label that is the target of a GOTO or some + other jump. Do not try to sort out the details, just reject + this loop nest. */ + { + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "Possible control transfer to label at %L " + "prevents annotation of loop nest", + &code->loc); + return ar_invalid_nest; + } + + switch (code->op) + { + case EXEC_OACC_KERNELS: + /* Enter kernels region. */ + annotate_do_loops_in_kernels (code->block->next, code, + goto_targets, + as_in_kernels_region); + walk_block = false; + break; + + case EXEC_OACC_PARALLEL_LOOP: + case EXEC_OACC_PARALLEL: + case EXEC_OACC_KERNELS_LOOP: + case EXEC_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. */ + if (state >= as_in_kernels_region) + { + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "Explicit loop annotation at %L " + "prevents annotation of loop nest", + &code->loc); + return ar_invalid_nest; + } + walk_block = false; + break; + + case EXEC_DO: + if (state >= as_in_kernels_region) + { + /* A DO loop's body is another phony DO node whose next + pointer starts the actual body. Skip the phony node. */ + gcc_assert (code->block->op == EXEC_DO); + enum annotation_result result + = annotate_do_loops_in_kernels (code->block->next, code, + goto_targets, + as_in_kernels_loop); + /* Check for function/procedure calls in the body of the + loop that would prevent parallelization. Unlike in C/C++, + we do not have to check that there is no modification of + the loop variable or loop count since they are already + handled by the semantics of DO loops in the FORTRAN + language. */ + if (result != ar_invalid_nest && check_for_invalid_calls (code)) + result = ar_invalid_nest; + if (result == ar_ok) + annotate_do_loop (code, parent); + else if (result == ar_invalid_nest + && state >= as_in_kernels_loop) + /* The outer loop is invalid, too, so stop traversal. */ + return result; + walk_block = false; + } + break; + + case EXEC_DO_WHILE: + case EXEC_DO_CONCURRENT: + /* Traverse the body in a special state to allow EXIT statements + from these loops. */ + if (state >= as_in_kernels_loop) + { + enum annotation_result result + = annotate_do_loops_in_kernels (code->block, code, + goto_targets, + as_in_kernels_inner_loop); + if (result == ar_invalid_nest) + return result; + else if (result != ar_ok) + retval = result; + walk_block = false; + } + break; + + case EXEC_GOTO: + case EXEC_ARITHMETIC_IF: + case EXEC_STOP: + case EXEC_ERROR_STOP: + /* A jump that may leave this loop. */ + if (state >= as_in_kernels_loop) + { + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "Possible unstructured control flow at %L " + "prevents annotation of loop nest", + &code->loc); + return ar_invalid_nest; + } + break; + + case EXEC_RETURN: + /* A return from a kernels region is diagnosed elsewhere as a + hard error, so no warning is needed here. */ + if (state >= as_in_kernels_loop) + return ar_invalid_nest; + break; + + case EXEC_EXIT: + if (state == as_in_kernels_loop) + { + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "Exit at %L prevents annotation of loop", + &code->loc); + retval = ar_invalid_loop; + } + break; + + case EXEC_BACKSPACE: + case EXEC_CLOSE: + case EXEC_ENDFILE: + case EXEC_FLUSH: + case EXEC_INQUIRE: + case EXEC_OPEN: + case EXEC_READ: + case EXEC_REWIND: + case EXEC_WRITE: + /* Executing side-effecting I/O statements in parallel doesn't + make much sense. If this is what users want, they can always + add explicit annotations on the loop nest. */ + if (state >= as_in_kernels_loop) + { + gfc_warning (OPT_Wopenacc_kernels_annotate_loops, + "I/O statement at %L prevents annotation of loop", + &code->loc); + return ar_invalid_nest; + } + break; + + default: + break; + } + + /* Visit nested statements, if any, returning early if we hit + any problems. */ + if (walk_block) + { + enum annotation_result result + = annotate_do_loops_in_kernels (code->block, code, + goto_targets, state); + if (result == ar_invalid_nest) + return result; + else if (result != ar_ok) + retval = result; + } + } + return retval; +} + +/* Traverse CODE to find all the labels referenced by GOTO and similar + statements and store them in GOTO_TARGETS. */ + +static void +compute_goto_targets (gfc_code *code, hash_set *goto_targets) +{ + for ( ; code; code = code->next) + { + switch (code->op) + { + case EXEC_GOTO: + case EXEC_LABEL_ASSIGN: + goto_targets->add (code->label1); + gcc_fallthrough (); + + case EXEC_ARITHMETIC_IF: + goto_targets->add (code->label2); + goto_targets->add (code->label3); + gcc_fallthrough (); + + default: + /* Visit nested statements, if any. */ + if (code->block != NULL) + compute_goto_targets (code->block, goto_targets); + } + } +} + +/* Find DO loops in OpenACC kernels regions that do not have OpenACC + annotations but look like they might benefit from automatic + parallelization. Add "acc loop auto" annotations for them. Assumes + flag_openacc_kernels_annotate_loops is set. */ + +void +gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *ns) +{ + if (ns->proc_name) + { + hash_set goto_targets; + compute_goto_targets (ns->code, &goto_targets); + annotate_do_loops_in_kernels (ns->code, NULL, &goto_targets, as_outer); + } + + for (ns = ns->contained; ns; ns = ns->sibling) + gfc_oacc_annotate_loops_in_kernels_regions (ns); +} diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c index 12aa80ec45ca..04e9d2450b16 100644 --- a/gcc/fortran/parse.c +++ b/gcc/fortran/parse.c @@ -6912,6 +6912,15 @@ done: if (flag_c_prototypes || flag_c_prototypes_external) fprintf (stdout, "\n#ifdef __cplusplus\n}\n#endif\n"); + /* Add annotations on loops in OpenACC kernels regions if requested. This + is most easily done on this representation close to the source code. */ + if (flag_openacc && flag_openacc_kernels_annotate_loops) + { + gfc_current_ns = gfc_global_ns_list; + for (; gfc_current_ns; gfc_current_ns = gfc_current_ns->sibling) + gfc_oacc_annotate_loops_in_kernels_regions (gfc_current_ns); + } + /* Do the translation. */ translate_all_program_units (gfc_global_ns_list); diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index 3fb48b321f2f..2ceae2088070 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -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/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index 6c8d298e236d..d061a241074b 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -2,6 +2,7 @@ ! 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/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 index 5defe2ea85de..d2816c3e9364 100644 --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 @@ -1,4 +1,5 @@ ! { dg-options "-fopenacc -fdump-tree-omplower" } +! { dg-additional-options "-fno-openacc-kernels-annotate-loops" } module consts integer, parameter :: n = 100 diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 index ef53324dd2a0..63774ffb5aff 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 @@ -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/gfortran.dg/goacc/kernels-loop-annotation-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95 new file mode 100644 index 000000000000..41f6307dbb17 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95 @@ -0,0 +1,33 @@ +! { 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. + +subroutine f (a, b, c) + implicit none + + real, intent (in), dimension(16,16) :: a + real, intent (in), dimension(16,16) :: b + real, intent (out), dimension(16,16) :: c + + integer :: i, j, k + real :: t + +!$acc kernels copyin(a(1:16,1:16), b(1:16,1:16)) copyout(c(1:16,1:16)) + + do i = 1, 16 + do j = 1, 16 + t = 0 + do k = 1, 16 + t = t + a(i,k) * b(k,j) + end do + c(i,j) = t; + end do + end do + +!$acc end kernels +end subroutine f + +! { dg-final { scan-tree-dump-times "acc loop private\\(.\\) auto" 3 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95 new file mode 100644 index 000000000000..f612c5beb963 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95 @@ -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. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + if (a(i) < 0 .or. b(i) < 0) then + go to 10 ! { dg-warning "Possible unstructured control flow" } + end if + t = t + a(i) * b(i) + end do + +10 f = t + +!$acc end kernels + +end function f diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95 new file mode 100644 index 000000000000..d51482e4685d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95 @@ -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-additional-options "-std=legacy" } +! { dg-do compile } + +! Test that a loop with a random label in the body cannot be annotated. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + t = 0.0 + +!$acc kernels + + goto 10 + + do i = 1, 16 +10 t = t + a(i) * b(i) ! { dg-warning "Possible control transfer to label" } + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95 new file mode 100644 index 000000000000..3c4956d70775 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95 @@ -0,0 +1,39 @@ +! { 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. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i, j + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + do j = 1, 16 + if (a(i) < 0 .or. b(j) < 0) then + exit ! { dg-warning "Exit" } + else + t = t + a(i) * b(j) + end if + end do + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95 new file mode 100644 index 000000000000..3ec459f0a8df --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95 @@ -0,0 +1,38 @@ +! { 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. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i, j + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + if (a(i) < 0) then + exit ! { dg-warning "Exit" } + end if + do j = 1, 16 + t = t + a(i) * b(j) + end do + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95 new file mode 100644 index 000000000000..91f431cca432 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95 @@ -0,0 +1,35 @@ +! { 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. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i, j + real :: t + + t = 0.0 + +!$acc kernels + +!$acc loop seq ! { dg-warning "Explicit loop annotation" } + do i = 1, 16 + do j = 1, 16 + t = t + a(i) * b(j) + end do + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95 new file mode 100644 index 000000000000..570c12d3ad70 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95 @@ -0,0 +1,35 @@ +! { 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 the outer loop, and produces a diagnostic. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i, j + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + !$acc loop seq ! { dg-warning "Explicit loop annotation" } + do j = 1, 16 + t = t + a(i) * b(j) + end do + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95 new file mode 100644 index 000000000000..6e44a304b28b --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95 @@ -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 loops containing I/O statements can't be annotated. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i, j + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + do j = 1, 16 + print *, " i =", i, " j =", j ! { dg-warning "I/O statement" } + t = t + a(i) * b(j) + end do + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95 new file mode 100644 index 000000000000..4624a05247d9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95 @@ -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 variable bound can be annotated. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (:) :: a, b + + integer :: i, n + real :: t + + t = 0.0 + n = size (a) + +!$acc kernels + + do i = 1, n + t = t + a(i) * b(i) + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95 new file mode 100644 index 000000000000..daed8f7f6e9d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95 @@ -0,0 +1,33 @@ +! { 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. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + if (a(i) > 0 .and. b(i) > 0) then + t = t + a(i) * b(i) + end if + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95 new file mode 100644 index 000000000000..0c4ad256b7eb --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95 @@ -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 case construct in the body can be annotated. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + +!$acc kernels + + do i = 1, 16 + select case (i) + case (1) + t = a(i) * b(i) + case default + t = t + a(i) * b(i) + end select + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95 new file mode 100644 index 000000000000..1c3f87eed6e4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95 @@ -0,0 +1,35 @@ +! { 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 cycle statement in the body can be annotated. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + if (a(i) < 0 .or. b(i) < 0) then + cycle + end if + t = t + a(i) * b(i) + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } } + diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95 new file mode 100644 index 000000000000..43173a70df24 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95 @@ -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 exit statement in the body cannot be annotated. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + if (a(i) < 0 .or. b(i) < 0) then + exit ! { dg-warning "Exit" } + end if + t = t + a(i) * b(i) + end do + + f = t + +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95 new file mode 100644 index 000000000000..ec42213220e7 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95 @@ -0,0 +1,48 @@ +! { 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. + + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + interface + function g (x) + real :: g + real, intent (in) :: x + end function g + + subroutine h (x) + real, intent (in) :: x + end subroutine h + end interface + + t = 0.0 + +!$acc kernels + do i = 1, 16 + t = t + g (a(i) * b(i)) ! { dg-warning "Function call" } + end do + + do i = 1, 16 + call h (t) ! { dg-warning "Subroutine call" } + t = t + a(i) * b(i) + end do + + f = t +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } } + diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95 new file mode 100644 index 000000000000..9188f70d9664 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95 @@ -0,0 +1,50 @@ +! { 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 call to a declared openacc function/subroutine +! can be annotated. + + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + interface + function g (x) + !$acc routine worker + real :: g + real, intent (in) :: x + end function g + + subroutine h (x) + !$acc routine worker + real, intent (in) :: x + end subroutine h + end interface + + t = 0.0 + +!$acc kernels + do i = 1, 16 + t = t + g (a(i) * b(i)) + end do + + do i = 1, 16 + call h (t) + t = t + a(i) * b(i) + end do + + f = t +!$acc end kernels + +end function f + +! { dg-final { scan-tree-dump-times "acc loop private\\(i\\) auto" 2 "original" } } + diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95 new file mode 100644 index 000000000000..f5aa5a0f43b5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95 @@ -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 return statement in the body gives a hard +! error. + +function f (a, b) + implicit none + + real :: f + real, intent (in), dimension (16) :: a, b + + integer :: i + real :: t + + t = 0.0 + +!$acc kernels + + do i = 1, 16 + if (a(i) < 0 .or. b(i) < 0) then + f = 0.0 + return ! { dg-error "invalid branch" } + end if + t = t + a(i) * b(i) + end do + + f = t + +!$acc end kernels + +end function f diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 index 2f1dcd603a14..c1f6ef8df600 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 @@ -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/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 index 447e85d64483..313e3df7f63d 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 @@ -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/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 index 4edb2889b7b1..26671064ba27 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 @@ -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/gfortran.dg/goacc/kernels-loop-data-update.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 index fc113e1f6602..d79ed796c366 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 @@ -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/gfortran.dg/goacc/kernels-loop-data.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 index 94522f586362..d8ef52af2e6a 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 @@ -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/gfortran.dg/goacc/kernels-loop-n.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95 index b9c4aea074d7..6b7334144c87 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95 @@ -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/gfortran.dg/goacc/kernels-loop.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 index 6dc7b2e0f28f..aadfcfc41448 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 @@ -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/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 index 48c20b999423..0d45c5cf4338 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 @@ -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" }