Kernels loops annotation: Fortran.

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

2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
	    Gergö Barany <gergo@codesourcery.com>

	gcc/fortran/
	* gfortran.h (gfc_oacc_annotate_loops_in_kernels_regions): Declare.
	* lang.opt (Wopenacc-kernels-annotate-loops): New.
	(fopenacc-kernels-annotate-loops): New.
	* openmp.cc: Include options.h.
	(enum annotation_state, enum annotation_result): New.
	(check_code_for_invalid_calls): New.
	(check_expr_for_invalid_calls): New.
	(check_for_invalid_calls): New.
	(annotate_do_loop): New.
	(annotate_do_loops_in_kernels): New.
	(compute_goto_targets): New.
	(gfc_oacc_annotate_loops_in_kernels_regions): New.
	* parse.cc (gfc_parse_file): Handle -fopenacc-kernels-annotate-loops.

	gcc/testsuite/
	* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add
	-fno-openacc-kernels-annotate-loops option.
	* gfortran.dg/goacc/classify-kernels.f95: Likewise.
	* gfortran.dg/goacc/common-block-3.f90: Likewise.
	* gfortran.dg/goacc/kernels-loop-2.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-data-2.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-data-update.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-data.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-n.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop.f95: Likewise.
	* gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95:
	Likewise.
	* gfortran.dg/goacc/kernels-loop-annotation-1.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-2.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-3.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-4.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-5.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-6.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-7.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-8.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-9.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-10.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-11.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-12.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-13.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-14.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-15.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-16.f95: New.

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 5dfe3858978ca9af6562da714e5f70d8ce91f2c6..a28bc3dc7892dc47703dcdbf00b1eb4d60ce35b9 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,20 @@
+2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
+           Gergö Barany <gergo@codesourcery.com>
+
+       * gfortran.h (gfc_oacc_annotate_loops_in_kernels_regions): Declare.
+       * lang.opt (Wopenacc-kernels-annotate-loops): New.
+       (fopenacc-kernels-annotate-loops): New.
+       * openmp.cc: Include options.h.
+       (enum annotation_state, enum annotation_result): New.
+       (check_code_for_invalid_calls): New.
+       (check_expr_for_invalid_calls): New.
+       (check_for_invalid_calls): New.
+       (annotate_do_loop): New.
+       (annotate_do_loops_in_kernels): New.
+       (compute_goto_targets): New.
+       (gfc_oacc_annotate_loops_in_kernels_regions): New.
+       * parse.cc (gfc_parse_file): Handle -fopenacc-kernels-annotate-loops.
+
 2020-02-19  Julian Brown  <julian@codesourcery.com>
 
        * trans-openmp.ccc (gfc_omp_check_optional_argument): Handle non-decl

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index b1d782702360ba1eed13bcb47eea9297c83f0166..b190f7a059a135f7c1ec82aae8298c5ead2befc3 100644 (file)
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -3581,6 +3581,7 @@ void gfc_resolve_oacc_directive (gfc_code *, gfc_namespace *);
 void gfc_resolve_oacc_declare (gfc_namespace *);
 void gfc_resolve_oacc_blocks (gfc_code *, gfc_namespace *);
 void gfc_resolve_oacc_routines (gfc_namespace *);
+void gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *);
 
 /* expr.cc */
 void gfc_free_actual_arglist (gfc_actual_arglist *);

diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt
index cf397129d55b528495c68d23d3b003b5b18167b5..c1b41b0006a104e7a0b71e765daa36dad9052420 100644 (file)
--- a/gcc/fortran/lang.opt
+++ b/gcc/fortran/lang.opt
@@ -289,6 +289,10 @@ Wopenacc-parallelism
 Fortran
 ; Documented in C
 
+Wopenacc-kernels-annotate-loops
+Fortran
+; Documented in C
+
 Wopenmp-simd
 Fortran
 ; Documented in C
@@ -704,6 +708,10 @@ fopenacc-dim=
 Fortran LTO Joined Var(flag_openacc_dims)
 ; Documented in C
 
+fopenacc-kernels-annotate-loops
+Fortran LTO Optimization
+; Documented in C
+
 fopenmp
 Fortran LTO
 ; Documented in C

diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 0e108c9474f0416f39a600a044b968a765574d8a..39a4e651e9519b3c9c99ddca9d38661a108cea4f 100644 (file)
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -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 '!'.  */
@@ -9464,3 +9465,366 @@ gfc_resolve_omp_udrs (gfc_symtree *st)
   for (omp_udr = st->n.omp_udr; omp_udr; omp_udr = omp_udr->next)
     gfc_resolve_omp_udr (omp_udr);
 }
+
+
+/* The following functions implement automatic recognition and annotation of
+   DO loops in OpenACC kernels regions.  Inside a kernels region, a nest of
+   DO loops that does not contain any annotated OpenACC loops, nor EXIT
+   or GOTO statements, gets an automatic "acc loop auto" annotation
+   on each loop.
+   This feature is controlled by flag_openacc_kernels_annotate_loops.  */
+
+
+/* State of annotation state traversal for DO loops in kernels regions.  */
+enum annotation_state {
+  as_outer,
+  as_in_kernels_region,
+  as_in_kernels_loop,
+  as_in_kernels_inner_loop
+};
+
+/* Return status of annotation traversal.  */
+enum annotation_result {
+  ar_ok,
+  ar_invalid_loop,
+  ar_invalid_nest
+};
+
+/* Code walk function for check_for_invalid_calls.  */
+
+static int
+check_code_for_invalid_calls (gfc_code **codep, int *walk_subtrees,
+                             void *data ATTRIBUTE_UNUSED)
+{
+  gfc_code *code = *codep;
+  switch (code->op)
+    {
+    case EXEC_CALL:
+      /* Calls to openacc routines are permitted.  */
+      if (code->resolved_sym
+         && (code->resolved_sym->attr.oacc_routine_lop
+             != OACC_ROUTINE_LOP_NONE))
+       return 0;
+      /* Else fall through.  */
+
+    case EXEC_CALL_PPC:
+    case EXEC_ASSIGN_CALL:
+      gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                  "Subroutine call at %L prevents annotation of loop nest",
+                  &code->loc);
+      *walk_subtrees = 0;
+      return 1;
+
+    default:
+      return 0;
+    }
+}
+
+/* Expr walk function for check_for_invalid_calls.  */
+
+static int
+check_expr_for_invalid_calls (gfc_expr **exprp, int *walk_subtrees,
+                             void *data ATTRIBUTE_UNUSED)
+{
+  gfc_expr *expr = *exprp;
+  switch (expr->expr_type)
+    {
+    case EXPR_FUNCTION:
+      if (expr->value.function.esym
+         && (expr->value.function.esym->attr.oacc_routine_lop
+             != OACC_ROUTINE_LOP_NONE))
+       return 0;
+      /* Else fall through.  */
+
+    case EXPR_COMPCALL:
+      gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                  "Function call at %L prevents annotation of loop nest",
+                  &expr->where);
+      *walk_subtrees = 0;
+      return 1;
+
+    default:
+      return 0;
+    }
+}
+
+/* Return TRUE if the DO loop CODE contains function or procedure
+   calls that ought to prohibit annotation.  This traversal is
+   separate from the main annotation tree walk because we need to walk
+   expressions as well as executable statements.  */
+
+static bool
+check_for_invalid_calls (gfc_code *code)
+{
+  gcc_assert (code->op == EXEC_DO);
+  return gfc_code_walker (&code, check_code_for_invalid_calls,
+                         check_expr_for_invalid_calls, NULL);
+}
+
+/* Annotate DO loop CODE with OpenACC "loop auto".  */
+
+static void
+annotate_do_loop (gfc_code *code, gfc_code *parent)
+{
+
+  /* A DO loop's body is another phony DO node whose next pointer starts
+     the actual body.  */
+  gcc_assert (code->op == EXEC_DO);
+  gcc_assert (code->block->op == EXEC_DO);
+
+  /* Build the "acc loop auto" annotation and add the loop as its
+     body.  */
+  gfc_omp_clauses *clauses = gfc_get_omp_clauses ();
+  clauses->par_auto = 1;
+  gfc_code *oacc_loop = gfc_get_code (EXEC_OACC_LOOP);
+  oacc_loop->block = gfc_get_code (EXEC_OACC_LOOP);
+  oacc_loop->block->next = code;
+  oacc_loop->ext.omp_clauses = clauses;
+  oacc_loop->loc = code->loc;
+  oacc_loop->block->loc = code->loc;
+
+  /* Splice the annotation into the place of the original loop.  */
+  if (parent->block == code)
+    parent->block = oacc_loop;
+  else
+    {
+      gfc_code *prev = parent->block;
+      while (prev != code && prev->next != code)
+       {
+         prev = prev->next;
+         gcc_assert (prev != NULL);
+       }
+      prev->next = oacc_loop;
+    }
+  oacc_loop->next = code->next;
+  code->next = NULL;
+}
+
+/* Recursively traverse CODE in block PARENT, finding OpenACC kernels
+   regions.  GOTO_TARGETS keeps track of statement labels that are
+   targets of gotos in the current function, while STATE keeps track
+   of the current context of the traversal.  If the traversal
+   encounters a DO loop inside a kernels region, annotate it with
+   OpenACC loop directives if appropriate.  Return the status of the
+   traversal.  */
+
+static enum annotation_result
+annotate_do_loops_in_kernels (gfc_code *code, gfc_code *parent,
+                             hash_set <gfc_st_label *> *goto_targets,
+                             annotation_state state)
+{
+  gfc_code *next_code = NULL;
+  enum annotation_result retval = ar_ok;
+
+  for ( ; code; code = next_code)
+    {
+      bool walk_block = true;
+      next_code = code->next;
+
+      if (state >= as_in_kernels_loop
+         && code->here && goto_targets->contains (code->here))
+       /* This statement has a label that is the target of a GOTO or some
+          other jump.  Do not try to sort out the details, just reject
+          this loop nest.  */
+       {
+         gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                      "Possible control transfer to label at %L "
+                      "prevents annotation of loop nest",
+                      &code->loc);
+         return ar_invalid_nest;
+       }
+
+      switch (code->op)
+       {
+       case EXEC_OACC_KERNELS:
+         /* Enter kernels region.  */
+         annotate_do_loops_in_kernels (code->block->next, code,
+                                       goto_targets,
+                                       as_in_kernels_region);
+         walk_block = false;
+         break;
+
+       case EXEC_OACC_PARALLEL_LOOP:
+       case EXEC_OACC_PARALLEL:
+       case EXEC_OACC_KERNELS_LOOP:
+       case EXEC_OACC_LOOP:
+         /* Do not try to add automatic OpenACC annotations inside manually
+            annotated loops.  Presumably, the user avoided doing it on
+            purpose; for example, all available levels of parallelism may
+            have been used up.  */
+         if (state >= as_in_kernels_region)
+           {
+             gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                          "Explicit loop annotation at %L "
+                          "prevents annotation of loop nest",
+                          &code->loc);
+             return ar_invalid_nest;
+           }
+         walk_block = false;
+         break;
+
+       case EXEC_DO:
+         if (state >= as_in_kernels_region)
+           {
+             /* A DO loop's body is another phony DO node whose next
+                pointer starts the actual body.  Skip the phony node.  */
+             gcc_assert (code->block->op == EXEC_DO);
+             enum annotation_result result
+               = annotate_do_loops_in_kernels (code->block->next, code,
+                                               goto_targets,
+                                               as_in_kernels_loop);
+             /* Check for function/procedure calls in the body of the
+                loop that would prevent parallelization.  Unlike in C/C++,
+                we do not have to check that there is no modification of
+                the loop variable or loop count since they are already
+                handled by the semantics of DO loops in the FORTRAN
+                language.  */
+             if (result != ar_invalid_nest && check_for_invalid_calls (code))
+               result = ar_invalid_nest;
+             if (result == ar_ok)
+               annotate_do_loop (code, parent);
+             else if (result == ar_invalid_nest
+                      && state >= as_in_kernels_loop)
+               /* The outer loop is invalid, too, so stop traversal.  */
+               return result;
+             walk_block = false;
+           }
+         break;
+
+       case EXEC_DO_WHILE:
+       case EXEC_DO_CONCURRENT:
+         /* Traverse the body in a special state to allow EXIT statements
+            from these loops.  */
+         if (state >= as_in_kernels_loop)
+           {
+             enum annotation_result result
+               = annotate_do_loops_in_kernels (code->block, code,
+                                               goto_targets,
+                                               as_in_kernels_inner_loop);
+             if (result == ar_invalid_nest)
+               return result;
+             else if (result != ar_ok)
+               retval = result;
+             walk_block = false;
+           }
+         break;
+
+       case EXEC_GOTO:
+       case EXEC_ARITHMETIC_IF:
+       case EXEC_STOP:
+       case EXEC_ERROR_STOP:
+         /* A jump that may leave this loop.  */
+         if (state >= as_in_kernels_loop)
+           {
+             gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                          "Possible unstructured control flow at %L "
+                          "prevents annotation of loop nest",
+                          &code->loc);
+             return ar_invalid_nest;
+           }
+         break;
+
+       case EXEC_RETURN:
+         /* A return from a kernels region is diagnosed elsewhere as a
+            hard error, so no warning is needed here.  */
+         if (state >= as_in_kernels_loop)
+           return ar_invalid_nest;
+         break;
+
+       case EXEC_EXIT:
+         if (state == as_in_kernels_loop)
+           {
+             gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                          "Exit at %L prevents annotation of loop",
+                          &code->loc);
+             retval = ar_invalid_loop;
+           }
+         break;
+
+       case EXEC_BACKSPACE:
+       case EXEC_CLOSE:
+       case EXEC_ENDFILE:
+       case EXEC_FLUSH:
+       case EXEC_INQUIRE:
+       case EXEC_OPEN:
+       case EXEC_READ:
+       case EXEC_REWIND:
+       case EXEC_WRITE:
+         /* Executing side-effecting I/O statements in parallel doesn't
+            make much sense.  If this is what users want, they can always
+            add explicit annotations on the loop nest.  */
+         if (state >= as_in_kernels_loop)
+           {
+             gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+                          "I/O statement at %L prevents annotation of loop",
+                          &code->loc);
+             return ar_invalid_nest;
+           }
+         break;
+
+       default:
+         break;
+       }
+
+      /* Visit nested statements, if any, returning early if we hit
+        any problems.  */
+      if (walk_block)
+       {
+         enum annotation_result result
+           = annotate_do_loops_in_kernels (code->block, code,
+                                           goto_targets, state);
+         if (result == ar_invalid_nest)
+           return result;
+         else if (result != ar_ok)
+           retval = result;
+       }
+    }
+  return retval;
+}
+
+/* Traverse CODE to find all the labels referenced by GOTO and similar
+   statements and store them in GOTO_TARGETS.  */
+
+static void
+compute_goto_targets (gfc_code *code, hash_set <gfc_st_label *> *goto_targets)
+{
+  for ( ; code; code = code->next)
+    {
+      switch (code->op)
+       {
+       case EXEC_GOTO:
+       case EXEC_LABEL_ASSIGN:
+         goto_targets->add (code->label1);
+         gcc_fallthrough ();
+
+       case EXEC_ARITHMETIC_IF:
+         goto_targets->add (code->label2);
+         goto_targets->add (code->label3);
+         gcc_fallthrough ();
+
+       default:
+         /* Visit nested statements, if any.  */
+         if (code->block != NULL)
+           compute_goto_targets (code->block, goto_targets);
+       }
+    }
+}
+
+/* Find DO loops in OpenACC kernels regions that do not have OpenACC
+   annotations but look like they might benefit from automatic
+   parallelization.  Add "acc loop auto" annotations for them.  Assumes
+   flag_openacc_kernels_annotate_loops is set.  */
+
+void
+gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *ns)
+{
+  if (ns->proc_name)
+    {
+      hash_set <gfc_st_label *> goto_targets;
+      compute_goto_targets (ns->code, &goto_targets);
+      annotate_do_loops_in_kernels (ns->code, NULL, &goto_targets, as_outer);
+    }
+
+  for (ns = ns->contained; ns; ns = ns->sibling)
+    gfc_oacc_annotate_loops_in_kernels_regions (ns);
+}

diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index e6e915d2a5e5afb4030ad4f6bdcdde3ecf47d0ea..bf38f70341da43b02c74670208846b3dd8d8b3fa 100644 (file)
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -6952,6 +6952,15 @@ done:
   if (flag_c_prototypes || flag_c_prototypes_external)
     fprintf (stdout, "\n#ifdef __cplusplus\n}\n#endif\n");
 
+  /* Add annotations on loops in OpenACC kernels regions if requested.  This
+     is most easily done on this representation close to the source code.  */
+  if (flag_openacc && flag_openacc_kernels_annotate_loops)
+    {
+      gfc_current_ns = gfc_global_ns_list;
+      for (; gfc_current_ns; gfc_current_ns = gfc_current_ns->sibling)
+       gfc_oacc_annotate_loops_in_kernels_regions (gfc_current_ns);
+    }
+
   /* Do the translation.  */
   translate_all_program_units (gfc_global_ns_list);
 

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index e59a826b9e2ea7ae430b5698ee76582f120bd8f0..44e9841c1c88709c2721cd1f8d87a79abb6be4c1 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,37 @@
+2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
+           Gergö Barany <gergo@codesourcery.com>
+
+       * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add
+       -fno-openacc-kernels-annotate-loops option.
+       * gfortran.dg/goacc/classify-kernels.f95: Likewise.
+       * gfortran.dg/goacc/common-block-3.f90: Likewise.
+       * gfortran.dg/goacc/kernels-loop-2.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop-data-2.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop-data-update.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop-data.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop-n.f95: Likewise.
+       * gfortran.dg/goacc/kernels-loop.f95: Likewise.
+       * gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95:
+       Likewise.
+       * gfortran.dg/goacc/kernels-loop-annotation-1.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-2.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-3.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-4.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-5.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-6.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-7.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-8.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-9.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-10.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-11.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-12.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-13.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-14.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-15.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-16.f95: New.
+
 2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
 
        * c-c++-common/goacc/kernels-decompose-2.c: Add

diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
index f8897fc5b3402a87603d4b038762f2389e1cd739..a867fe96bb9907fe6d08b4528739057dbddaac24 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
@@ -4,6 +4,7 @@
 ! { dg-additional-options "--param openacc-kernels=decompose" }
 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fopt-info-all-omp" }
 ! { dg-additional-options "-fdump-tree-ompexp" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
index 2ed6cdb6115726ce15904d89cb7439379fc24fe9..62653e2890f7fe0afa2c4bff3e0bb7a45252c329 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
@@ -4,6 +4,7 @@
 ! { dg-additional-options "--param openacc-kernels=decompose" }
 
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fopt-info-all-omp" }
 ! { dg-additional-options "-fdump-tree-ompexp" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }

diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index 6f08d7eb8d5b60ae29bba0b4a39d9164596b132d..0c151547e328ffad92fe1c9a2cc883069c6ab6d9 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -1,4 +1,5 @@
 ! { dg-options "-fopenacc -fdump-tree-omplower" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 
 ! { dg-additional-options "--param=openacc-kernels=decompose" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
index 831218604d8749a7d63e822ddc092e8356e7d697..3884a81e6fa32ae1065c29060afebefde7c8237d 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-2.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95
new file mode 100644 (file)
index 0000000..41f6307
--- /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 (file)
index 0000000..f612c5b
--- /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 (file)
index 0000000..d51482e
--- /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 (file)
index 0000000..3c4956d
--- /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 (file)
index 0000000..3ec459f
--- /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 (file)
index 0000000..91f431c
--- /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 (file)
index 0000000..570c12d
--- /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 (file)
index 0000000..6e44a30
--- /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 (file)
index 0000000..4624a05
--- /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 (file)
index 0000000..daed8f7
--- /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 (file)
index 0000000..0c4ad25
--- /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 (file)
index 0000000..1c3f87e
--- /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 (file)
index 0000000..43173a7
--- /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 (file)
index 0000000..ec42213
--- /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 (file)
index 0000000..9188f70
--- /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 (file)
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 357056ea96f2decd1c9a3239e60964c022756077..62cd4d38a36c32a75478439b21e975dda41bb120 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-2.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
index 01674e3e0889134698b65ee7656ea37a43df4eab..461279d75726e37044a9ca1c426b34c34a09ba2a 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
index ea690a8933f996bfee7450619216dd3edb1ee58e..cfb3803501b8e6bb54bedabe54ecc43027aa8529 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-enter-exit.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
index 64b5e66aab9efe135ff5e525aee76708dda6fd92..a9ec77b4db60e24398da307bf66faf4e9b7e47e6 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data-update.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
index e81af3cc2fa6f94471a8b05474c2226267522646..53a082b782e6204dccdc59cd96980e62d4ef40dc 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-data.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
index b899bd65cdaeb2cdb483cb3ccf5bbd521e42d893..0d1021e30235e83f4ddeee56d70f5961196336a4 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-n.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
index 563229de75bdfa7eca7dbe81cf4df4521d691b19..2f8db4b9fd7b3ab9d9f85ec6c6b9af203faebe39 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
index 4479ae37791ddca6456519c8c3b4c3af5c242cf8..7e9589de799ce7bb97ac76ae6e757d206470f605 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95
@@ -1,6 +1,7 @@
 ! { dg-additional-options "--param=openacc-kernels=parloops" } as this is
 ! specifically testing "parloops" handling.
 ! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
 ! { dg-additional-options "-fdump-tree-parloops1-all" }
 ! { dg-additional-options "-fdump-tree-optimized" }