[03/40] Kernels loops annotation: Fortran.

Message ID 20211215155447.19379-4-frederik@codesourcery.com
State New
Headers
Series OpenACC "kernels" Improvements |

Commit Message

Frederik Harwath Dec. 15, 2021, 3:54 p.m. UTC
  From: Sandra Loosemore <sandra@codesourcery.com>

This patch implements the Fortran support for adding "#pragma acc loop auto"
annotations to loops in OpenACC kernels regions.  It implements the same
-fopenacc-kernels-annotate-loops and -Wopenacc-kernels-annotate-loops options
that were previously added (and documented) for the C/C++ front ends.

Co-Authored-By: Gergö Barany <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.c: Include options.h.
        (enum annotation_state, enum annotation_result): New.
        (check_code_for_invalid_calls): New.
        (check_expr_for_invalid_calls): New.
        (check_for_invalid_calls): New.
        (annotate_do_loop): New.
        (annotate_do_loops_in_kernels): New.
        (compute_goto_targets): New.
        (gfc_oacc_annotate_loops_in_kernels_regions): New.
        * parse.c (gfc_parse_file): Handle -fopenacc-kernels-annotate-loops.

gcc/testsuite/
        * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add
        -fno-openacc-kernels-annotate-loops option.
        * gfortran.dg/goacc/classify-kernels.f95: Likewise.
        * gfortran.dg/goacc/common-block-3.f90: Likewise.
        * gfortran.dg/goacc/kernels-loop-2.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop-data-2.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop-data-update.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop-data.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop-n.f95: Likewise.
        * gfortran.dg/goacc/kernels-loop.f95: Likewise.
        * gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95:
        Likewise.
        * gfortran.dg/goacc/kernels-loop-annotation-1.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-2.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-3.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-4.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-5.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-6.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-7.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-8.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-9.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-10.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-11.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-12.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-13.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-14.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-15.f95: New.
        * gfortran.dg/goacc/kernels-loop-annotation-16.f95: New.
---
 gcc/fortran/gfortran.h                        |   1 +
 gcc/fortran/lang.opt                          |   8 +
 gcc/fortran/openmp.c                          | 364 ++++++++++++++++++
 gcc/fortran/parse.c                           |   9 +
 .../goacc/classify-kernels-unparallelized.f95 |   1 +
 .../gfortran.dg/goacc/classify-kernels.f95    |   1 +
 .../gfortran.dg/goacc/common-block-3.f90      |   1 +
 .../gfortran.dg/goacc/kernels-loop-2.f95      |   1 +
 .../goacc/kernels-loop-annotation-1.f95       |  33 ++
 .../goacc/kernels-loop-annotation-10.f95      |  32 ++
 .../goacc/kernels-loop-annotation-11.f95      |  34 ++
 .../goacc/kernels-loop-annotation-12.f95      |  39 ++
 .../goacc/kernels-loop-annotation-13.f95      |  38 ++
 .../goacc/kernels-loop-annotation-14.f95      |  35 ++
 .../goacc/kernels-loop-annotation-15.f95      |  35 ++
 .../goacc/kernels-loop-annotation-16.f95      |  34 ++
 .../goacc/kernels-loop-annotation-2.f95       |  32 ++
 .../goacc/kernels-loop-annotation-3.f95       |  33 ++
 .../goacc/kernels-loop-annotation-4.f95       |  34 ++
 .../goacc/kernels-loop-annotation-5.f95       |  35 ++
 .../goacc/kernels-loop-annotation-6.f95       |  34 ++
 .../goacc/kernels-loop-annotation-7.f95       |  48 +++
 .../goacc/kernels-loop-annotation-8.f95       |  50 +++
 .../goacc/kernels-loop-annotation-9.f95       |  34 ++
 .../gfortran.dg/goacc/kernels-loop-data-2.f95 |   1 +
 .../goacc/kernels-loop-data-enter-exit-2.f95  |   1 +
 .../goacc/kernels-loop-data-enter-exit.f95    |   1 +
 .../goacc/kernels-loop-data-update.f95        |   1 +
 .../gfortran.dg/goacc/kernels-loop-data.f95   |   1 +
 .../gfortran.dg/goacc/kernels-loop-n.f95      |   1 +
 .../gfortran.dg/goacc/kernels-loop.f95        |   1 +
 .../kernels-parallel-loop-data-enter-exit.f95 |   1 +
 32 files changed, 974 insertions(+)
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95

--
2.33.0

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Patch

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index f7662c59a5df..50db768ce0fc 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -3545,6 +3545,7 @@  void gfc_resolve_oacc_declare (gfc_namespace *);
 void gfc_resolve_oacc_parallel_loop_blocks (gfc_code *, gfc_namespace *);
 void gfc_resolve_oacc_blocks (gfc_code *, gfc_namespace *);
 void gfc_resolve_oacc_routines (gfc_namespace *);
+void gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *);

 /* expr.c */
 void gfc_free_actual_arglist (gfc_actual_arglist *);
diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt
index 6db01c736be1..a202c04c4a25 100644
--- a/gcc/fortran/lang.opt
+++ b/gcc/fortran/lang.opt
@@ -289,6 +289,10 @@  Wopenacc-parallelism
 Fortran
 ; Documented in C

+Wopenacc-kernels-annotate-loops
+Fortran
+; Documented in C
+
 Wopenmp-simd
 Fortran
 ; Documented in C
@@ -695,6 +699,10 @@  fopenacc-dim=
 Fortran LTO Joined Var(flag_openacc_dims)
 ; Documented in C

+fopenacc-kernels-annotate-loops
+Fortran LTO Optimization
+; Documented in C
+
 fopenmp
 Fortran LTO
 ; Documented in C
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index dcf22ac2c2f3..243b5e0a9ac6 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -29,6 +29,7 @@  along with GCC; see the file COPYING3.  If not see
 #include "diagnostic.h"
 #include "gomp-constants.h"
 #include "target-memory.h"  /* For gfc_encode_character.  */
+#include "options.h"

 /* Match an end of OpenMP directive.  End of OpenMP directive is optional
    whitespace, followed by '\n' or comment '!'.  */
@@ -9090,3 +9091,366 @@  gfc_resolve_omp_udrs (gfc_symtree *st)
   for (omp_udr = st->n.omp_udr; omp_udr; omp_udr = omp_udr->next)
     gfc_resolve_omp_udr (omp_udr);
 }
+
+
+/* The following functions implement automatic recognition and annotation of
+   DO loops in OpenACC kernels regions.  Inside a kernels region, a nest of
+   DO loops that does not contain any annotated OpenACC loops, nor EXIT
+   or GOTO statements, gets an automatic "acc loop auto" annotation
+   on each loop.
+   This feature is controlled by flag_openacc_kernels_annotate_loops.  */
+
+
+/* State of annotation state traversal for DO loops in kernels regions.  */
+enum annotation_state {
+  as_outer,
+  as_in_kernels_region,
+  as_in_kernels_loop,
+  as_in_kernels_inner_loop
+};
+
+/* Return status of annotation traversal.  */
+enum annotation_result {
+  ar_ok,
+  ar_invalid_loop,
+  ar_invalid_nest
+};
+
+/* Code walk function for check_for_invalid_calls.  */
+
+static int
+check_code_for_invalid_calls (gfc_code **codep, int *walk_subtrees,
+                             void *data ATTRIBUTE_UNUSED)
+{
+  gfc_code *code = *codep;
+  switch (code->op)
+    {
+    case EXEC_CALL:
+      /* Calls to openacc routines are permitted.  */
+      if (code->resolved_sym
+         && (code->resolved_sym->attr.oacc_routine_lop
+             != OACC_ROUTINE_LOP_NONE))
+       return 0;
+      /* Else fall through.  */
+
+    case EXEC_CALL_PPC:
+    case EXEC_ASSIGN_CALL:
+      gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                  "Subroutine call at %L prevents annotation of loop nest",
+                  &code->loc);
+      *walk_subtrees = 0;
+      return 1;
+
+    default:
+      return 0;
+    }
+}
+
+/* Expr walk function for check_for_invalid_calls.  */
+
+static int
+check_expr_for_invalid_calls (gfc_expr **exprp, int *walk_subtrees,
+                             void *data ATTRIBUTE_UNUSED)
+{
+  gfc_expr *expr = *exprp;
+  switch (expr->expr_type)
+    {
+    case EXPR_FUNCTION:
+      if (expr->value.function.esym
+         && (expr->value.function.esym->attr.oacc_routine_lop
+             != OACC_ROUTINE_LOP_NONE))
+       return 0;
+      /* Else fall through.  */
+
+    case EXPR_COMPCALL:
+      gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                  "Function call at %L prevents annotation of loop nest",
+                  &expr->where);
+      *walk_subtrees = 0;
+      return 1;
+
+    default:
+      return 0;
+    }
+}
+
+/* Return TRUE if the DO loop CODE contains function or procedure
+   calls that ought to prohibit annotation.  This traversal is
+   separate from the main annotation tree walk because we need to walk
+   expressions as well as executable statements.  */
+
+static bool
+check_for_invalid_calls (gfc_code *code)
+{
+  gcc_assert (code->op == EXEC_DO);
+  return gfc_code_walker (&code, check_code_for_invalid_calls,
+                         check_expr_for_invalid_calls, NULL);
+}
+
+/* Annotate DO loop CODE with OpenACC "loop auto".  */
+
+static void
+annotate_do_loop (gfc_code *code, gfc_code *parent)
+{
+
+  /* A DO loop's body is another phony DO node whose next pointer starts
+     the actual body.  */
+  gcc_assert (code->op == EXEC_DO);
+  gcc_assert (code->block->op == EXEC_DO);
+
+  /* Build the "acc loop auto" annotation and add the loop as its
+     body.  */
+  gfc_omp_clauses *clauses = gfc_get_omp_clauses ();
+  clauses->par_auto = 1;
+  gfc_code *oacc_loop = gfc_get_code (EXEC_OACC_LOOP);
+  oacc_loop->block = gfc_get_code (EXEC_OACC_LOOP);
+  oacc_loop->block->next = code;
+  oacc_loop->ext.omp_clauses = clauses;
+  oacc_loop->loc = code->loc;
+  oacc_loop->block->loc = code->loc;
+
+  /* Splice the annotation into the place of the original loop.  */
+  if (parent->block == code)
+    parent->block = oacc_loop;
+  else
+    {
+      gfc_code *prev = parent->block;
+      while (prev != code && prev->next != code)
+       {
+         prev = prev->next;
+         gcc_assert (prev != NULL);
+       }
+      prev->next = oacc_loop;
+    }
+  oacc_loop->next = code->next;
+  code->next = NULL;
+}
+
+/* Recursively traverse CODE in block PARENT, finding OpenACC kernels
+   regions.  GOTO_TARGETS keeps track of statement labels that are
+   targets of gotos in the current function, while STATE keeps track
+   of the current context of the traversal.  If the traversal
+   encounters a DO loop inside a kernels region, annotate it with
+   OpenACC loop directives if appropriate.  Return the status of the
+   traversal.  */
+
+static enum annotation_result
+annotate_do_loops_in_kernels (gfc_code *code, gfc_code *parent,
+                             hash_set <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.c b/gcc/fortran/parse.c
index 12aa80ec45ca..04e9d2450b16 100644
--- a/gcc/fortran/parse.c
+++ b/gcc/fortran/parse.c
@@ -6912,6 +6912,15 @@  done:
   if (flag_c_prototypes || flag_c_prototypes_external)
     fprintf (stdout, "\n#ifdef __cplusplus\n}\n#endif\n");

+  /* Add annotations on loops in OpenACC kernels regions if requested.  This
+     is most easily done on this representation close to the source code.  */
+  if (flag_openacc && flag_openacc_kernels_annotate_loops)
+    {
+      gfc_current_ns = gfc_global_ns_list;
+      for (; gfc_current_ns; gfc_current_ns = gfc_current_ns->sibling)
+       gfc_oacc_annotate_loops_in_kernels_regions (gfc_current_ns);
+    }
+
   /* Do the translation.  */
   translate_all_program_units (gfc_global_ns_list);

diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
index 3fb48b321f2f..2ceae2088070 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
@@ -2,6 +2,7 @@ 
 ! OpenACC kernels.

 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fopt-info-optimized-omp" }
 ! { dg-additional-options "-fdump-tree-ompexp" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
index 6c8d298e236d..d061a241074b 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
@@ -2,6 +2,7 @@ 
 ! kernels.

 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fopt-info-optimized-omp" }
 ! { dg-additional-options "-fdump-tree-ompexp" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index 5defe2ea85de..d2816c3e9364 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -1,4 +1,5 @@ 
 ! { dg-options "-fopenacc -fdump-tree-omplower" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }

 module consts
   integer, parameter :: n = 100
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
index ef53324dd2a0..63774ffb5aff 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95
new file mode 100644
index 000000000000..41f6307dbb17
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95
@@ -0,0 +1,33 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that all loops in the nest are annotated.
+
+subroutine f (a, b, c)
+  implicit none
+
+  real, intent (in), dimension(16,16) :: a
+  real, intent (in), dimension(16,16) :: b
+  real, intent (out), dimension(16,16) :: c
+
+  integer :: i, j, k
+  real :: t
+
+!$acc kernels copyin(a(1:16,1:16), b(1:16,1:16)) copyout(c(1:16,1:16))
+
+  do i = 1, 16
+    do j = 1, 16
+      t = 0
+      do k = 1, 16
+        t = t + a(i,k) * b(k,j)
+      end do
+      c(i,j) = t;
+    end do
+  end do
+
+!$acc end kernels
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop private\\(.\\) auto" 3 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95
new file mode 100644
index 000000000000..f612c5beb963
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95
@@ -0,0 +1,32 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a random goto in the body can't be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    if (a(i) < 0 .or. b(i) < 0) then
+      go to 10  ! { dg-warning "Possible unstructured control flow" }
+    end if
+    t = t + a(i) * b(i)
+  end do
+
+10  f = t
+
+!$acc end kernels
+
+end function f
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95
new file mode 100644
index 000000000000..d51482e4685d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95
@@ -0,0 +1,34 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-std=legacy" }
+! { dg-do compile }
+
+! Test that a loop with a random label in the body cannot be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  goto 10
+
+  do i = 1, 16
+10  t = t + a(i) * b(i)  ! { dg-warning "Possible control transfer to label" }
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95
new file mode 100644
index 000000000000..3c4956d70775
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95
@@ -0,0 +1,39 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that in a situation with nested loops, a problem that prevents
+! annotation of the inner loop only still allows the outer loop to be
+! annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i, j
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    do j = 1, 16
+      if (a(i) < 0 .or. b(j) < 0) then
+        exit  ! { dg-warning "Exit" }
+      else
+        t = t + a(i) * b(j)
+      end if
+    end do
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95
new file mode 100644
index 000000000000..3ec459f0a8df
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95
@@ -0,0 +1,38 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that in a situation with nested loops, a problem that prevents
+! annotation of the outer loop only still allows the inner loop to be
+! annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i, j
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    if (a(i) < 0) then
+      exit  ! { dg-warning "Exit" }
+    end if
+    do j = 1, 16
+      t = t + a(i) * b(j)
+    end do
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95
new file mode 100644
index 000000000000..91f431cca432
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95
@@ -0,0 +1,35 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that an explicit annotation on an outer loop suppresses annotation
+!  of inner loops, and produces a diagnostic.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i, j
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+!$acc loop seq  ! { dg-warning "Explicit loop annotation" }
+  do i = 1, 16
+    do j = 1, 16
+      t = t + a(i) * b(j)
+    end do
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95
new file mode 100644
index 000000000000..570c12d3ad70
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95
@@ -0,0 +1,35 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that an explicit annotation on an inner loop suppresses annotation
+! of the outer loop, and produces a diagnostic.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i, j
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    !$acc loop seq  ! { dg-warning "Explicit loop annotation" }
+    do j = 1, 16
+      t = t + a(i) * b(j)
+    end do
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95
new file mode 100644
index 000000000000..6e44a304b28b
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95
@@ -0,0 +1,34 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that loops containing I/O statements can't be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i, j
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    do j = 1, 16
+      print *, " i =", i, " j =", j  ! { dg-warning "I/O statement" }
+      t = t + a(i) * b(j)
+    end do
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95
new file mode 100644
index 000000000000..4624a05247d9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95
@@ -0,0 +1,32 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a variable bound can be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (:) :: a, b
+
+  integer :: i, n
+  real :: t
+
+  t = 0.0
+  n = size (a)
+
+!$acc kernels
+
+  do i = 1, n
+    t = t + a(i) * b(i)
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95
new file mode 100644
index 000000000000..daed8f7f6e9d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95
@@ -0,0 +1,33 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a conditional in the body can be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    if (a(i) > 0 .and. b(i) > 0) then
+      t = t + a(i) * b(i)
+    end if
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95
new file mode 100644
index 000000000000..0c4ad256b7eb
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95
@@ -0,0 +1,34 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a case construct in the body can be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+!$acc kernels
+
+  do i = 1, 16
+    select case (i)
+      case (1)
+        t = a(i) * b(i)
+      case default
+        t = t + a(i) * b(i)
+    end select
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95
new file mode 100644
index 000000000000..1c3f87eed6e4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95
@@ -0,0 +1,35 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a cycle statement in the body can be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    if (a(i) < 0 .or. b(i) < 0) then
+      cycle
+    end if
+    t = t + a(i) * b(i)
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95
new file mode 100644
index 000000000000..43173a70df24
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95
@@ -0,0 +1,34 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a exit statement in the body cannot be annotated.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    if (a(i) < 0 .or. b(i) < 0) then
+      exit     ! { dg-warning "Exit" }
+    end if
+    t = t + a(i) * b(i)
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95
new file mode 100644
index 000000000000..ec42213220e7
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95
@@ -0,0 +1,48 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a random function call in the body cannot
+! be annotated.
+
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  interface
+    function g (x)
+      real :: g
+      real, intent (in) :: x
+    end function g
+
+    subroutine h (x)
+      real, intent (in) :: x
+    end subroutine h
+  end interface
+
+  t = 0.0
+
+!$acc kernels
+  do i = 1, 16
+    t = t + g (a(i) * b(i))  ! { dg-warning "Function call" }
+  end do
+
+  do i = 1, 16
+    call h (t) ! { dg-warning "Subroutine call" }
+    t = t + a(i) * b(i)
+  end do
+
+  f = t
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95
new file mode 100644
index 000000000000..9188f70d9664
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95
@@ -0,0 +1,50 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a call to a declared openacc function/subroutine
+! can be annotated.
+
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  interface
+    function g (x)
+      !$acc routine worker
+      real :: g
+      real, intent (in) :: x
+    end function g
+
+    subroutine h (x)
+      !$acc routine worker
+      real, intent (in) :: x
+    end subroutine h
+  end interface
+
+  t = 0.0
+
+!$acc kernels
+  do i = 1, 16
+    t = t + g (a(i) * b(i))
+  end do
+
+  do i = 1, 16
+    call h (t)
+    t = t + a(i) * b(i)
+  end do
+
+  f = t
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private\\(i\\) auto" 2 "original" } }
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95
new file mode 100644
index 000000000000..f5aa5a0f43b5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95
@@ -0,0 +1,34 @@ 
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a return statement in the body gives a hard
+! error.
+
+function f (a, b)
+  implicit none
+
+  real :: f
+  real, intent (in), dimension (16) :: a, b
+
+  integer :: i
+  real :: t
+
+  t = 0.0
+
+!$acc kernels
+
+  do i = 1, 16
+    if (a(i) < 0 .or. b(i) < 0) then
+      f = 0.0
+      return   ! { dg-error "invalid branch" }
+    end if
+    t = t + a(i) * b(i)
+  end do
+
+  f = t
+
+!$acc end kernels
+
+end function f
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
index 2f1dcd603a14..c1f6ef8df600 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
index 447e85d64483..313e3df7f63d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
index 4edb2889b7b1..26671064ba27 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
index fc113e1f6602..d79ed796c366 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
index 94522f586362..d8ef52af2e6a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
index b9c4aea074d7..6b7334144c87 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
index 6dc7b2e0f28f..aadfcfc41448 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
index 48c20b999423..0d45c5cf4338 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
@@ -1,4 +1,5 @@ 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }