[gcc/devel/omp/gcc-11] Relax some restrictions on the loop bound in kernels loop annotation.

[gcc/devel/omp/gcc-11] Relax some restrictions on the loop bound in kernels loop annotation.

[Kwok Yeung]

https://gcc.gnu.org/g:37108578754992c4a3a14e89c84921508d439500

commit 37108578754992c4a3a14e89c84921508d439500
Author: Sandra Loosemore <sandra@codesourcery.com>
Date:   Sun Aug 30 12:15:23 2020 -0700

    Relax some restrictions on the loop bound in kernels loop annotation.
    
    OpenACC loop semantics require that the loop bound be computable
    before entering the loop, rather than the C/C++ semantics where the
    end test is evaluated on every iteration.  Formerly the kernels loop
    annotater permitted only constants and variables not modified in the
    loop body in the loop bound expression.  This patch relaxes those
    restrictions somewhat to allow many forms of expressions involving
    such constants and variables, including calls to constant functions.
    
    2020-08-30  Sandra Loosemore  <sandra@codesourcery.com>
    
            gcc/c-family/
            * c-omp.c (end_test_ok_for_annotation_r): New.
            (end_test_ok_for_annotation): New.
            (check_and_annotate_for_loop): Use the new helper function.
    
            gcc/testsuite/
            * c-c++-common/goacc/kernels-loop-annotation-21.c: New.
            * c-c++-common/goacc/kernels-loop-annotation-22.c: New.

Diff:
---
 gcc/c-family/ChangeLog.omp                         |   6 ++
 gcc/c-family/c-omp.c                               | 120 +++++++++++++++++++--
 gcc/testsuite/ChangeLog.omp                        |   5 +
 .../goacc/kernels-loop-annotation-21.c             |  42 ++++++++
 .../goacc/kernels-loop-annotation-22.c             |  41 +++++++
 5 files changed, 205 insertions(+), 9 deletions(-)

diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index daf7c3ba596..34570313969 100644
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,3 +1,9 @@
+2020-08-30  Sandra Loosemore  <sandra@codesourcery.com>
+
+	* c-omp.c (end_test_ok_for_annotation_r): New.
+	(end_test_ok_for_annotation): New.
+	(check_and_annotate_for_loop): Use the new helper function.
+
 2020-08-30  Sandra Loosemore  <sandra@codesourcery.com>
     
 	* c-omp.c (annotate_for_loop): Move initializer processing...
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index ae2599c7859..2460b2af5a2 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2937,6 +2937,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
@@ -3105,16 +3215,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index f0f48bc978f..ac2bd2df8bb 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2020-08-30  Sandra Loosemore  <sandra@codesourcery.com>
+
+	* c-c++-common/goacc/kernels-loop-annotation-21.c: New.
+	* c-c++-common/goacc/kernels-loop-annotation-22.c: New.
+
 2020-08-24  Tobias Burnus  <tobias@codesourcery.com>
 	    Sandra Loosemore  <sandra@codesourcery.com>
 
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 00000000000..f87444ede4b
--- /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 00000000000..6a5099d2ff9
--- /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" } } */