[12/40] Relax some restrictions on the loop bound in kernels loop annotation.

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index e73fb5d01f7e..dc63d304ca67 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3165,6 +3165,116 @@  is_local_var (tree decl)
          && !TREE_ADDRESSABLE (decl));
 }

+/* EXP is a loop bound expression for a comparison against local
+   variable DECL.  Check whether this is potentially valid in an OpenACC loop
+   context, namely that it can be precomputed when entering the loop
+   construct per the OpenACC specification.  Local variables referenced
+   in both DECL and EXP that may not be modified in the body of the loop
+   are added to the list in INFO to be checked later.
+
+   FIXME: Ideally we would like to make this test permissive rather than
+   restrictive, and allow the later conversion of the "auto" attribute to
+   either "seq" or "independent" to make the determination using dataflow,
+   alias analysis, etc rather than a tree traversal.  But presently it does
+   not do that and always just hoists the loop bound expression.  So the
+   current implementation only considers expressions involving unmodified
+   local variables and constants, using a tree walk.  */
+
+static tree
+end_test_ok_for_annotation_r (tree *tp, int *walk_subtrees,
+                             void *data)
+{
+  tree exp = *tp;
+  struct annotation_info *info = (struct annotation_info *) data;
+
+  switch (TREE_CODE_CLASS (TREE_CODE (exp)))
+    {
+    case tcc_constant:
+      /* Constants are trivially known to be invariant.  */
+      return NULL_TREE;
+
+    case tcc_declaration:
+      if (is_local_var (exp))
+       {
+         tree t;
+         /* Add it to the list of variables that can't be modified in the
+            loop, only if not already present.  */
+         for (t = info->vars; t && TREE_VALUE (t) != exp;
+              t = TREE_CHAIN (t))
+           ;
+         if (!t)
+           info->vars = tree_cons (NULL_TREE, exp, info->vars);
+         return NULL_TREE;
+       }
+      else if (TREE_CODE (exp) == VAR_DECL && TREE_READONLY (exp))
+       return NULL_TREE;
+      else if (TREE_CODE (exp) == FUNCTION_DECL)
+       return NULL_TREE;
+      break;
+
+    case tcc_unary:
+    case tcc_binary:
+    case tcc_comparison:
+      /* Allow arithmetic expressions and comparisons provided
+        that the operands are good.  */
+      return NULL_TREE;
+
+    default:
+      /* Handle some special cases.  */
+      switch (TREE_CODE (exp))
+       {
+       case COND_EXPR:
+       case TRUTH_ANDIF_EXPR:
+       case TRUTH_ORIF_EXPR:
+       case TRUTH_AND_EXPR:
+       case TRUTH_OR_EXPR:
+       case TRUTH_XOR_EXPR:
+       case TRUTH_NOT_EXPR:
+         /* ?: and boolean operators are OK.  */
+         return NULL_TREE;
+
+       case CALL_EXPR:
+         /* Allow calls to constant functions with invariant operands.  */
+         {
+           tree fndecl = get_callee_fndecl (exp);
+           if (fndecl && TREE_READONLY (fndecl))
+             return NULL_TREE;
+         }
+         break;
+
+       case ADDR_EXPR:
+         /* We can expect addresses of things to be invariant.  */
+         return NULL_TREE;
+
+       default:
+         break;
+       }
+    }
+
+  /* Reject anything else.  */
+  *walk_subtrees = 0;
+  return exp;
+}
+
+static bool
+end_test_ok_for_annotation (tree decl, tree exp,
+                           struct annotation_info *info)
+{
+  /* Traversal returns NULL_TREE if all is well.  */
+  if (!walk_tree (&exp, end_test_ok_for_annotation_r, info, NULL))
+    {
+      /* So far, so good.  Check the decl against any variables collected
+        in the exp.  */
+      tree t;
+      for (t = info->vars; t; t = TREE_CHAIN (t))
+       if (TREE_VALUE (t) == decl)
+         return false;
+      info->vars = tree_cons (NULL_TREE, decl, info->vars);
+      return true;
+    }
+  return false;
+}
+
 /* The initializer for a FOR_STMT is sometimes wrapped in various other
    language-specific tree structures.  We need a hook to unwrap them.
    This function takes a tree argument and should return either a
@@ -3333,16 +3443,8 @@  check_and_annotate_for_loop (tree *nodeptr, tree_stmt_iterator *prev_tsi,
        limit_exp = TREE_OPERAND (cond, 0);

       if (!limit_exp
-         || (!is_local_var (limit_exp)
-             && (TREE_CODE_CLASS (TREE_CODE (limit_exp)) != tcc_constant)))
+         || !end_test_ok_for_annotation (decl, limit_exp, &loop_info))
        do_not_annotate_loop (&loop_info, as_invalid_predicate, cond);
-      else
-       {
-         /* These variables must not be assigned to in the loop.  */
-         loop_info.vars = tree_cons (NULL_TREE, decl, loop_info.vars);
-         if (TREE_CODE_CLASS (TREE_CODE (limit_exp)) != tcc_constant)
-           loop_info.vars = tree_cons (NULL_TREE, limit_exp, loop_info.vars);
-       }
     }

   /* Walk the body.  This will process any nested loops, so we have to do it
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c
new file mode 100644
index 000000000000..f87444ede4b4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-21.c
@@ -0,0 +1,42 @@ 
+/* { 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 for rejecting annotation on loops that have various subexpressions
+   in the loop end test that are not loop-invariant.  */
+
+extern int g (int);
+extern int x;
+extern int gg (int, int) __attribute__ ((const));
+
+void f (float *a, float *b, int n)
+{
+
+  int j;
+#pragma acc kernels
+  {
+    /* Non-constant function call.  */
+    for (int i = 0; i < g(n); i++)     /* { dg-warning "loop cannot be annotated" } */
+      a[i] = b[i];
+
+    /* Global variable.  */
+    for (int i = x; i < n + x; i++)    /* { dg-warning "loop cannot be annotated" } */
+      a[i] = b[i];
+
+    /* Explicit reference to the loop variable.  */
+    for (int i = 0; i < gg (i, n); i++)        /* { dg-warning "loop cannot be annotated" } */
+      a[i] = b[i];
+
+    /* Reference to a variable that is modified in the body of the loop.  */
+    j = 0;
+    for (int i = 0; i < gg (j, n); i++)        /* { dg-warning "loop cannot be annotated" } */
+      {
+       a[i] = b[i];
+       j = i;
+      }
+
+  }
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 0 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c
new file mode 100644
index 000000000000..6a5099d2ff9d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-22.c
@@ -0,0 +1,41 @@ 
+/* { 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 for accepting annotation on loops that have various forms of
+   loop-invariant expressions in their end test.  */
+
+extern const int x;
+extern int g (int) __attribute__ ((const));
+
+void f (float *a, float *b, int n)
+{
+
+  int j;
+#pragma acc kernels
+  {
+    /* Reversed form of comparison.  */
+    for (int i = 0; n >= i; i++)
+      a[i] = b[i];
+
+    /* Constant function call.  */
+    for (int i = 0; i < g(n); i++)
+      a[i] = b[i];
+
+    /* Constant global variable.  */
+    for (int i = 0; i < x; i++)
+      a[i] = b[i];
+
+    /* Complicated expression involving conditionals, etc. */
+    for (int i = 0; i < ((x == 4) ? (n << 2) : (n << 3)); i++)
+      a[i] = b[i];
+
+    /* Reference to a local variable not modified in the loop.  */
+    j = ((x == 4) ? (n << 2) : (n << 3));
+    for (int i = 0; i < j; i++)
+      a[i] = b[i];
+  }
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 5 "original" } } */