[gcc/devel/omp/gcc-13] Kernels loops annotation: Fortran.
Kwok Yeung
kcy@gcc.gnu.org
Fri May 19 16:40:28 GMT 2023
https://gcc.gnu.org/g:6bc5f50de4feaa34cf21c2a85e77ac742d7cfcf0
commit 6bc5f50de4feaa34cf21c2a85e77ac742d7cfcf0
Author: Sandra Loosemore <sandra@codesourcery.com>
Date: Mon Mar 16 18:08:01 2020 -0700
Kernels loops annotation: Fortran.
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.
2020-03-27 Sandra Loosemore <sandra@codesourcery.com>
Gergö Barany <gergo@codesourcery.com>
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.cc: 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.cc (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.
Diff:
---
gcc/fortran/ChangeLog.omp | 17 +
gcc/fortran/gfortran.h | 1 +
gcc/fortran/lang.opt | 8 +
gcc/fortran/openmp.cc | 364 +++++++++++++++++++++
gcc/fortran/parse.cc | 9 +
gcc/testsuite/ChangeLog.omp | 34 ++
.../goacc/classify-kernels-unparallelized.f95 | 1 +
.../gfortran.dg/goacc/classify-kernels.f95 | 1 +
gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 1 +
gcc/testsuite/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 +
.../gfortran.dg/goacc/kernels-loop-data-update.f95 | 1 +
.../gfortran.dg/goacc/kernels-loop-data.f95 | 1 +
gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95 | 1 +
gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 | 1 +
.../kernels-parallel-loop-data-enter-exit.f95 | 1 +
34 files changed, 1025 insertions(+)
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 5dfe3858978..a28bc3dc789 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,20 @@
+2020-03-27 Sandra Loosemore <sandra@codesourcery.com>
+ Gergö Barany <gergo@codesourcery.com>
+
+ * gfortran.h (gfc_oacc_annotate_loops_in_kernels_regions): Declare.
+ * lang.opt (Wopenacc-kernels-annotate-loops): New.
+ (fopenacc-kernels-annotate-loops): New.
+ * openmp.cc: 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.cc (gfc_parse_file): Handle -fopenacc-kernels-annotate-loops.
+
2020-02-19 Julian Brown <julian@codesourcery.com>
* trans-openmp.ccc (gfc_omp_check_optional_argument): Handle non-decl
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 8f83aeaccb6..772e04fe308 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -3630,6 +3630,7 @@ void gfc_resolve_oacc_directive (gfc_code *, gfc_namespace *);
void gfc_resolve_oacc_declare (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.cc */
void gfc_free_actual_arglist (gfc_actual_arglist *);
diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt
index 7236351a93c..1cf5e2c26c1 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
@@ -704,6 +708,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.cc b/gcc/fortran/openmp.cc
index 4e01869999a..387ef866557 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -30,6 +30,7 @@ along with GCC; see the file COPYING3. If not see
#include "gomp-constants.h"
#include "target-memory.h" /* For gfc_encode_character. */
#include "bitmap.h"
+#include "options.h"
static gfc_statement omp_code_to_statement (gfc_code *);
@@ -10380,3 +10381,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 <gfc_st_label *> *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 <gfc_st_label *> *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 <gfc_st_label *> 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.cc b/gcc/fortran/parse.cc
index f1e55316e5b..114a3ff5f10 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -7022,6 +7022,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index db031acbba1..7e4e695115d 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,37 @@
+2020-03-27 Sandra Loosemore <sandra@codesourcery.com>
+ Gergö Barany <gergo@codesourcery.com>
+
+ * 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.
+
2020-03-27 Sandra Loosemore <sandra@codesourcery.com>
* c-c++-common/goacc/kernels-decompose-2.c: Add
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
index e8ceda9b877..df975950013 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
@@ -4,6 +4,7 @@
! { dg-additional-options "--param openacc-kernels=decompose" }
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fopt-info-all-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 7eb79188b82..568822f67ce 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
@@ -4,6 +4,7 @@
! { dg-additional-options "--param openacc-kernels=decompose" }
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fopt-info-all-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 6f08d7eb8d5..0c151547e32 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" }
! { dg-additional-options "--param=openacc-kernels=decompose" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
index a5705393fa5..c7252697a7e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
@@ -1,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 00000000000..41f6307dbb1
--- /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 00000000000..f612c5beb96
--- /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 00000000000..d51482e4685
--- /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 00000000000..3c4956d7077
--- /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 00000000000..3ec459f0a8d
--- /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 00000000000..91f431cca43
--- /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 00000000000..570c12d3ad7
--- /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 00000000000..6e44a304b28
--- /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 00000000000..4624a05247d
--- /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 00000000000..daed8f7f6e9
--- /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 00000000000..0c4ad256b7e
--- /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 00000000000..1c3f87eed6e
--- /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 00000000000..43173a70df2
--- /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 00000000000..ec42213220e
--- /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 00000000000..9188f70d966
--- /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 00000000000..f5aa5a0f43b
--- /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 881126c1d81..714e8641abc 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
@@ -1,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 f90b7157ade..d748b58b965 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,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 9fe10f238e3..2afbf6497b8 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,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 d81ae60935f..869bc2f5ea6 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
@@ -1,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 7dcb5287983..c123699d27a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
@@ -1,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 227045950d3..78cfa7d79e8 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
@@ -1,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 29471df1935..7e1c7cb863c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
@@ -1,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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 dac6dc6eaeb..552bb618cbc 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,6 +1,7 @@
! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
! specifically testing "parloops" handling.
! { 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" }
More information about the Gcc-cvs
mailing list