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 8f83aeaccb6ce512edf73d5eb6a1c8b4ca9def86..772e04fe3087742871c94ba420c6ca91928690f4 100644 (file)
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -3630,6 +3630,7 @@ void gfc_resolve_oacc_directive (gfc_code *, gfc_namespace *);
 void gfc_resolve_oacc_declare (gfc_namespace *);
 void gfc_resolve_oacc_blocks (gfc_code *, gfc_namespace *);
 void gfc_resolve_oacc_routines (gfc_namespace *);
+void gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *);
 
 /* expr.cc */
 void gfc_free_actual_arglist (gfc_actual_arglist *);

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

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

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index db031acbba1e533780e83df2b7f5dba81c82c54d..7e4e695115dbe9a7b5099291ed1791eee094dafe 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 e8ceda9b877cdb02218cd8ba391346c4dcf20be1..df9759500134e710f39dc8ba52ecf832229d20cb 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 7eb79188b8269f1db989cf5def549c484e146f25..568822f67ce812412375455ba8ba51fa2eb7afc7 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 a5705393fa5dcff1198e96509e45045a37c17489..c7252697a7e1ee57d964e06288e994d273176d77 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 881126c1d8143e9782f5a9cf341871013d0e0bdc..714e8641abcc6a08e532906e72aea4f38a6b4581 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 f90b7157ade080dd69bc8519232b95c301c72d45..d748b58b965341ceac99fc1a78923b433e7142a3 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 9fe10f238e335ea6d0e0825135890e243aad8bc1..2afbf6497b8455c3d36df849e7e40d9c1d4206c6 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 d81ae60935f62d5740fc9862641f1cae5a58b0b3..869bc2f5ea6d648b08eaa956a0aae9bff119f365 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 7dcb5287983f7b31b6f7801a464a7785b1ab313c..c123699d27aef74f386559b8bb01207f3b5b9bf8 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 227045950d355da2043720587d02eb6bd1c7c0de..78cfa7d79e80546229eec62afc7957c5a984106f 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 29471df193583a452ae7e67e5872bafcce1be021..7e1c7cb863c75f402451967fd2e427f43b2dcb09 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 dac6dc6eaeb9761915f52766fda0e4a76477d847..552bb618cbc96d8f8404ecf44bd68129ae29a12e 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" }