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-09-08 Sandra Loosemore <san...@codesourcery.com>
Gergö Barany <ge...@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.c: Include options.h.
(enum annotation_state): New.
(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.
* trans-openmp.c (gfc_trans_omp_do): Add combined parameter.
Use it to set OACC_LOOP_COMBINED. Adjust call sites.
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/combined-directives.f90: Adjust patterns.
* gfortran.dg/goacc/common-block-3.f90: Add
-fno-openacc-kernels-annotate-loops option.
* 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/private-explicit-kernels-1.f95: Adjust
patterns.
* gfortran.dg/goacc/private-predetermined-kernels-1.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.
* gfortran.dg/goacc/kernels-loop-annotation-18.f95: New.
* gfortran.dg/goacc/kernels-loop-annotation-19.f95: New.
* gfortran.dg/goacc/kernels-loop-annotation-20.f95: New.
---
gcc/fortran/gfortran.h | 1 +
gcc/fortran/lang.opt | 8 +
gcc/fortran/openmp.c | 415 +++++++++++++++++++++
gcc/fortran/parse.c | 9 +
gcc/fortran/trans-openmp.c | 30 +-
.../goacc/classify-kernels-unparallelized.f95 | 1 +
.../gfortran.dg/goacc/classify-kernels.f95 | 1 +
.../gfortran.dg/goacc/combined-directives.f90 | 19 +-
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-18.f95 | 28 ++
.../goacc/kernels-loop-annotation-19.f95 | 29 ++
.../goacc/kernels-loop-annotation-2.f95 | 32 ++
.../goacc/kernels-loop-annotation-20.f95 | 26 ++
.../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 +
.../goacc/private-explicit-kernels-1.f95 | 7 +-
.../goacc/private-predetermined-kernels-1.f95 | 7 +-
39 files changed, 1152 insertions(+), 19 deletions(-)
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-18.f95
create mode 100644
gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.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-20.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
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index d0cea83..bbde046 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -3328,6 +3328,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 da4b1aa..34dc9ee 100644
--- a/gcc/fortran/lang.opt
+++ b/gcc/fortran/lang.opt
@@ -285,6 +285,10 @@ Wuse-without-only
Fortran Var(warn_use_without_only) Warning
Warn about USE statements that have no ONLY qualifier.
+Wopenacc-kernels-annotate-loops
+Fortran
+; Documented in C
+
Wopenmp-simd
Fortran
; Documented in C
@@ -687,6 +691,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 d0e516c..a7c4331 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see
#include "parse.h"
#include "diagnostic.h"
#include "gomp-constants.h"
+#include "options.h"
/* Match an end of OpenMP directive. End of OpenMP directive is optional
whitespace, followed by '\n' or comment '!'. */
@@ -7039,3 +7040,417 @@ 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:
+ /* Permit calls to Fortran intrinsic functions and to routines
+ with an explicitly declared parallelism level. */
+ if (expr->value.function.isym
+ || (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_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_OACC_KERNELS_LOOP:
+ /* This is a combined "acc kernels loop" directive. We want to
+ leave the outer loop alone but try to annotate any nested
+ loops in the body. The expected structure nesting here is
+ EXEC_OACC_KERNELS_LOOP
+ EXEC_OACC_KERNELS_LOOP
+ EXEC_DO
+ EXEC_DO
+ ...body... */
+ if (code->block)
+ /* Might be empty? */
+ {
+ gcc_assert (code->block->op == EXEC_OACC_KERNELS_LOOP);
+ gfc_omp_clauses *clauses = code->ext.omp_clauses;
+ int collapse = clauses->collapse;
+ gfc_expr_list *tile = clauses->tile_list;
+ gfc_code *inner = code->block->next;
+
+ gcc_assert (inner->op == EXEC_DO);
+ gcc_assert (inner->block->op == EXEC_DO);
+
+ /* We need to skip over nested loops covered by "collapse" or
+ "tile" clauses. "Tile" takes precedence
+ (see gfc_trans_omp_do). */
+ if (tile)
+ {
+ collapse = 0;
+ for (gfc_expr_list *el = tile; el; el = el->next)
+ collapse++;
+ }
+ if (clauses->orderedc)
+ collapse = clauses->orderedc;
+ if (collapse <= 0)
+ collapse = 1;
+ for (int i = 1; i < collapse; i++)
+ {
+ gcc_assert (inner->op == EXEC_DO);
+ gcc_assert (inner->block->op == EXEC_DO);
+ inner = inner->block->next;
+ }
+ if (inner)
+ /* Loop might have empty body? */
+ annotate_do_loops_in_kernels (inner->block->next,
+ inner, goto_targets,
+ as_in_kernels_region);
+ }
+ 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.c b/gcc/fortran/parse.c
index 6669621..a19286a 100644
--- a/gcc/fortran/parse.c
+++ b/gcc/fortran/parse.c
@@ -6612,6 +6612,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/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 7d3365f..8956530 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -4315,7 +4315,8 @@ typedef struct dovar_init_d {
static tree
gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock,
- gfc_omp_clauses *do_clauses, tree par_clauses)
+ gfc_omp_clauses *do_clauses, tree par_clauses,
+ bool combined)
{
gfc_se se;
tree dovar, stmt, from, to, step, type, init, cond, incr, orig_decls;
@@ -4645,7 +4646,10 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op,
stmtblock_t *pblock,
case EXEC_OMP_DO: stmt = make_node (OMP_FOR); break;
case EXEC_OMP_DISTRIBUTE: stmt = make_node (OMP_DISTRIBUTE); break;
case EXEC_OMP_TASKLOOP: stmt = make_node (OMP_TASKLOOP); break;
- case EXEC_OACC_LOOP: stmt = make_node (OACC_LOOP); break;
+ case EXEC_OACC_LOOP:
+ stmt = make_node (OACC_LOOP);
+ OACC_LOOP_COMBINED (stmt) = combined;
+ break;
default: gcc_unreachable ();
}
@@ -4739,7 +4743,8 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
pblock = █
else
pushlevel ();
- stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL);
+ stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL,
+ true);
protected_set_expr_location (stmt, loc);
if (TREE_CODE (stmt) != BIND_EXPR)
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
@@ -5180,7 +5185,7 @@ gfc_trans_omp_do_simd (gfc_code *code, stmtblock_t
*pblock,
omp_do_clauses
= gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DO], code->loc);
body = gfc_trans_omp_do (code, EXEC_OMP_SIMD, pblock ? pblock : &block,
- &clausesa[GFC_OMP_SPLIT_SIMD], omp_clauses);
+ &clausesa[GFC_OMP_SPLIT_SIMD], omp_clauses, false);
if (pblock == NULL)
{
if (TREE_CODE (body) != BIND_EXPR)
@@ -5233,7 +5238,7 @@ gfc_trans_omp_parallel_do (gfc_code *code, stmtblock_t
*pblock,
pushlevel ();
}
stmt = gfc_trans_omp_do (code, EXEC_OMP_DO, new_pblock,
- &clausesa[GFC_OMP_SPLIT_DO], omp_clauses);
+ &clausesa[GFC_OMP_SPLIT_DO], omp_clauses, false);
if (pblock == NULL)
{
if (TREE_CODE (stmt) != BIND_EXPR)
@@ -5476,7 +5481,8 @@ gfc_trans_omp_distribute (gfc_code *code, gfc_omp_clauses
*clausesa)
case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block,
- &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE);
+ &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE,
+ false);
if (TREE_CODE (stmt) != BIND_EXPR)
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
else
@@ -5532,7 +5538,7 @@ gfc_trans_omp_teams (gfc_code *code, gfc_omp_clauses
*clausesa,
case EXEC_OMP_TEAMS_DISTRIBUTE:
stmt = gfc_trans_omp_do (code, EXEC_OMP_DISTRIBUTE, NULL,
&clausesa[GFC_OMP_SPLIT_DISTRIBUTE],
- NULL);
+ NULL, false);
break;
default:
stmt = gfc_trans_omp_distribute (code, clausesa);
@@ -5600,7 +5606,8 @@ gfc_trans_omp_target (gfc_code *code)
break;
case EXEC_OMP_TARGET_SIMD:
stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block,
- &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE);
+ &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE,
+ false);
if (TREE_CODE (stmt) != BIND_EXPR)
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
else
@@ -5670,7 +5677,8 @@ gfc_trans_omp_taskloop (gfc_code *code)
break;
case EXEC_OMP_TASKLOOP_SIMD:
stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block,
- &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE);
+ &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE,
+ false);
if (TREE_CODE (stmt) != BIND_EXPR)
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
else
@@ -5948,7 +5956,7 @@ gfc_trans_oacc_directive (gfc_code *code)
return gfc_trans_oacc_construct (code);
case EXEC_OACC_LOOP:
return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
- NULL);
+ NULL, false);
case EXEC_OACC_UPDATE:
case EXEC_OACC_CACHE:
case EXEC_OACC_ENTER_DATA:
@@ -5985,7 +5993,7 @@ gfc_trans_omp_directive (gfc_code *code)
case EXEC_OMP_SIMD:
case EXEC_OMP_TASKLOOP:
return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
- NULL);
+ NULL, false);
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
case EXEC_OMP_DISTRIBUTE_SIMD:
diff --git
a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
index 0877242..fe8d09f 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 f2c4736..1a0b8b8 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/combined-directives.f90
b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
index 9563492..562a4e4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -139,10 +139,21 @@ end subroutine test
! { dg-final { scan-tree-dump-times "acc loop private.i. private.j.
collapse.2." 2 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop private.i. gang" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 2
"gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2
"gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2
"gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2
"gimple" } }
+
+! These are the parallel loop variants.
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 1
"gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 1
"gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 1
"gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 1
"gimple" } }
+
+! These are the kernels loop variants. Here the inner loops are annotated
+! separately.
+! { dg-final { scan-tree-dump-times "acc loop private.i. worker" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. vector" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. seq" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. auto" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop auto private.j." 4 "gimple" } }
+
! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3"
2 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2
"gimple" } }
! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index 5defe2e..d2816c3 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 ef53324..63774ff 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 0000000..42e751d
--- /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 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 0000000..f612c5be
--- /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 0000000..6e2e2c4
--- /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 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 0000000..03c4234
--- /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 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 0000000..6aeb3f2
--- /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 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 0000000..7d1cff6
--- /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 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 0000000..dab0d40
--- /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 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 0000000..15ef670
--- /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 auto" 0 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95
b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95
new file mode 100644
index 0000000..e4e210a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95
@@ -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 "acc kernels loop" directive causes annotation of the entire
+! loop nest.
+
+subroutine f (a, b)
+
+ implicit none
+ real, intent (in), dimension(20) :: a
+ real, intent (out), dimension(20) :: b
+ integer :: k, l, m
+
+!$acc kernels loop
+ do k = 1, 20
+ do l = 1, 20
+ do m = 1, 20
+ b(m) = a(m);
+ end do
+ end do
+ end do
+
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop auto" 2 "original" } }
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95
b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95
new file mode 100644
index 0000000..5dd6e7f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95
@@ -0,0 +1,29 @@
+! { 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 "acc kernels loop" directive causes annotation of the entire
+! loop nest in the presence of a collapse clause.
+
+subroutine f (a, b)
+
+ implicit none
+ real, intent (in), dimension(20) :: a
+ real, intent (out), dimension(20) :: b
+ integer :: k, l, m
+
+!$acc kernels loop collapse(2)
+ do k = 1, 20
+ do l = 1, 20
+ do m = 1, 20
+ b(m) = a(m);
+ end do
+ end do
+ end do
+
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop .*collapse.2." 1 "original" } }
+! { dg-final { scan-tree-dump-times "acc loop auto" 1 "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 0000000..2baaa59
--- /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 auto" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95
b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95
new file mode 100644
index 0000000..5169a0a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95
@@ -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 calls to intrinsics in the body can be annotated.
+
+subroutine f (n, input, out1, out2)
+ implicit none
+ integer :: n
+ integer, intent (in), dimension (n) :: input
+ integer, intent (out), dimension (n) :: out1, out2
+
+ integer :: i
+
+!$acc kernels
+
+ do i = 1, n
+ out1(i) = min (i, input(i))
+ out2(i) = not (input(i))
+ end do
+!$acc end kernels
+
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop 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 0000000..e629891
--- /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 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 0000000..6c3300b
--- /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 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 0000000..52a9e7e
--- /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 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 0000000..60eb245
--- /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 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 0000000..438a13a
--- /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 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 0000000..aa97e37
--- /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 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 0000000..f5aa5a0
--- /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 2f1dcd6..c1f6ef8 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 447e85d6..313e3df 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 4edb288..2667106 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 fc113e1..d79ed79 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 94522f5..d8ef52a 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 b9c4aea..6b73341 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 6dc7b2e..aadfcfc 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 48c20b9..0d45c5c 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" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
index 5d563d2..0c47045 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
@@ -73,8 +73,9 @@ program test
!$acc kernels loop private(i2_1_c, j2_1_c) independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\)
private\\(j2_1_c\\) independent" 1 "original" } }
- ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\)
private\\(j2_1_c\\) independent" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\)
independent" 1 "gimple" } }
do i2_1_c = 1, 100
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop auto
private\\(j2_1_c\\)" 1 "gimple" } }
do j2_1_c = 1, 100
end do
end do
@@ -130,9 +131,11 @@ program test
!$acc kernels loop private(i3_1_c, j3_1_c, k3_1_c) independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\)
private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
- ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\)
private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\)
independent" 1 "gimple" } }
do i3_1_c = 1, 100
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop auto
private\\(j3_1_c\\)" 1 "gimple" } }
do j3_1_c = 1, 100
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop auto
private\\(k3_1_c\\)" 1 "gimple" } }
do k3_1_c = 1, 100
end do
end do
diff --git
a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
index 12a7854..3357a20 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -73,8 +73,9 @@ program test
!$acc kernels loop independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\)
private\\(j2_1_c\\) independent" 1 "original" } }
- ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\)
private\\(j2_1_c\\) independent" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\)
independent" 1 "gimple" } }
do i2_1_c = 1, 100
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop auto
private\\(j2_1_c\\)" 1 "gimple" } }
do j2_1_c = 1, 100
end do
end do
@@ -130,9 +131,11 @@ program test
!$acc kernels loop independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\)
private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
- ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\)
private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\)
independent" 1 "gimple" } }
do i3_1_c = 1, 100
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop auto
private\\(j3_1_c\\)" 1 "gimple" } }
do j3_1_c = 1, 100
+ ! { dg-final { scan-tree-dump-times "#pragma acc loop auto
private\\(k3_1_c\\)" 1 "gimple" } }
do k3_1_c = 1, 100
end do
end do
--
2.8.1
