[PATCH 11/40] Clean up loop variable extraction in OpenACC kernels loop annotation.

[Frederik Harwath <frederik@codesourcery.com>]

From: Sandra Loosemore <sandra@codesourcery.com>

The code for identifying annotatable loops in OpenACC kernels regions
previously looked for the loop variable as the left-hand side of the
comparison in the loop end test.  However, front end optimizations
sometimes switch the sense of the comparison, making this method
unreliable.  In particular, it's ambiguous when both operands to the
end test comparison are local variables.

This patch reorders the loop processing to identify the loop variable
from the initializer, rather than the end test. The processing of the
end test then just checks that one of the operands to the comparison
matches the variable appearing in the initializer.  Much of the patch
is code refactoring, moving the initializer analysis out of
annotate_for_loop to check_and_annotate_for_loop so it can be
performed earlier.

2020-08-30  Sandra Loosemore  <sandra@codesourcery.com>

        gcc/c-family/
        * c-omp.c (annotate_for_loop): Move initializer processing...
        (check_and_annotate_for_loop): ... to here.  Allow the loop
        variable as either operand to the condition.
---
 gcc/c-family/c-omp.c | 196 +++++++++++++++++++++----------------------
 1 file changed, 98 insertions(+), 98 deletions(-)

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index e7c27f45e888..e73fb5d01f7e 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3174,86 +3174,26 @@ static tree (*lang_specific_unwrap_initializer) (tree);

 /* Try to annotate the given NODE, which must be a FOR_STMT, with a
    "#pragma acc loop auto" annotation.  In practice, this means
-   building an OMP_FOR node for it.  PREV_STMT is the statement
-   immediately before the loop, which may be used as the loop's
-   initialization statement.  Annotating the loop may fail, in which
-   case INFO is used to record the cause of the failure and the
-   original loop remains unchanged.  This function returns the
-   transformed loop if the transformation succeeded, the original node
-   otherwise.  */
+   building an OMP_FOR node for it.  DECL and INIT are the
+   previously-verified iteration variable and initializer.  Annotating
+   the loop may fail, in which case INFO is used to record the cause
+   of the failure and the original loop remains unchanged.  This
+   function returns the transformed loop if the transformation
+   succeeded, the original node otherwise.  */

 static tree
-annotate_for_loop (tree node, tree_stmt_iterator *prev_tsi,
+annotate_for_loop (tree node, tree decl, tree init,
                   struct annotation_info *info)
 {
   gcc_checking_assert (TREE_CODE (node) == FOR_STMT);

   location_t loc = EXPR_LOCATION (node);
   tree cond = FOR_COND (node);
+  tree incr = FOR_EXPR (node);
+
+  gcc_assert (decl);
   gcc_assert (cond);
-  tree decl = TREE_OPERAND (cond, 0);
   gcc_assert (decl && TREE_CODE (decl) == VAR_DECL);
-  tree init = FOR_INIT_STMT (node);
-  tree prev_stmt = NULL_TREE;
-  bool unlink_prev = false;
-  bool fix_decl = false;
-
-
-  /* Both the C and C++ front ends normally put the initializer in the
-     statement list just before the FOR_STMT instead of in FOR_INIT_STMT.
-     If FOR_INIT_STMT happens to exist but isn't a MODIFY_EXPR, bail out
-     because the code below won't handle it.  */
-  if (init != NULL_TREE && TREE_CODE (init) != MODIFY_EXPR)
-    {
-      do_not_annotate_loop (info, as_invalid_initializer, NULL_TREE);
-      return node;
-    }
-
-  /* Examine the statement before the loop to see if it is a
-     valid initializer.  It must be either a MODIFY_EXPR or VAR_DECL,
-     possibly wrapped in language-specific structure.  */
-  if (init == NULL_TREE && prev_tsi != NULL)
-    {
-      prev_stmt = tsi_stmt (*prev_tsi);
-
-      /* Call the language-specific hook to unwrap prev_stmt.  */
-      if (prev_stmt)
-       prev_stmt = (*lang_specific_unwrap_initializer) (prev_stmt);
-
-      /* See if we have a valid MODIFY_EXPR.  */
-      if (prev_stmt
-         && TREE_CODE (prev_stmt) == MODIFY_EXPR
-         && TREE_OPERAND (prev_stmt, 0) == decl
-         && !TREE_SIDE_EFFECTS (TREE_OPERAND (prev_stmt, 1)))
-       {
-         init = prev_stmt;
-         unlink_prev = true;
-       }
-      else if (prev_stmt == decl
-              && !TREE_SIDE_EFFECTS (DECL_INITIAL (decl)))
-       {
-         /* If the preceding statement is the declaration of the loop
-            variable with its initialization, build an assignment
-            expression for the loop's initializer.  */
-         init = build2 (MODIFY_EXPR, TREE_TYPE (decl), decl,
-                        DECL_INITIAL (decl));
-         /* We need to remove the initializer from the decl if we
-            end up using the init we just built instead.  */
-         fix_decl = true;
-       }
-    }
-
-  if (init == NULL_TREE)
-    /* There is nothing we can do to find the correct init statement for
-       this loop, but c_finish_omp_for insists on having one and would fail
-       otherwise.  In that case, we would just return node.  Do that
-       directly, here.  */
-    {
-      do_not_annotate_loop (info, as_missing_initializer, NULL_TREE);
-      return node;
-    }
-
-  tree incr = FOR_EXPR (node);

   /* The C++ frontend can wrap the increment two levels deep inside a
      cleanup expression, but c_finish_omp_for does not care about that.  */
@@ -3278,18 +3218,6 @@ annotate_for_loop (tree node, tree_stmt_iterator *prev_tsi,
                                            NULL_TREE, false, info);
   if (omp_for != NULL_TREE)
     {
-      if (unlink_prev)
-       /* We don't need the previous statement that we consumed as an
-          initializer in the new OMP_FOR any more.  */
-       tsi_delink (prev_tsi);
-
-      if (fix_decl)
-       /* We no longer need the initializer expression on the decl of
-          the loop variable and don't want to duplicate it.  The
-          kernels conversion pass would interpret it as a stray
-          assignment in a gang-single region.  */
-       DECL_INITIAL (prev_stmt) = NULL_TREE;
-
       /* Add an auto clause, then return the new loop.  */
       tree auto_clause = build_omp_clause (loc, OMP_CLAUSE_AUTO);
       OMP_CLAUSE_CHAIN (auto_clause) = OMP_FOR_CLAUSES (omp_for);
@@ -3315,11 +3243,16 @@ check_and_annotate_for_loop (tree *nodeptr, tree_stmt_iterator *prev_tsi,
 {
   tree node = *nodeptr;
   gcc_assert (TREE_CODE (node) == FOR_STMT);
+  tree init = FOR_INIT_STMT (node);
+  tree cond = FOR_COND (node);
+  tree prev_stmt = NULL_TREE;
+  tree decl = NULL_TREE;
+  bool unlink_prev = false;
+  bool fix_decl = false;

   /* This structure describes the current loop statement.  */
   struct annotation_info loop_info
     = { node, NULL_TREE, false, as_in_kernels_loop, NULL_TREE, info };
-  tree cond = FOR_COND (node);

   /* If we are in the body of an explicitly-annotated loop, do not add
      annotations to this loop or any other nested loops.  */
@@ -3331,30 +3264,84 @@ check_and_annotate_for_loop (tree *nodeptr, tree_stmt_iterator *prev_tsi,
      That is why we are doing some checks on the loop condition
      that duplicate what c_finish_omp_for is doing.  */

-  /* The loop condition must be a comparison.  */
+  /* First we need to find the decl and initializer for the
+     controlling variable.  Both the C and C++ front ends normally put
+     the initializer in the statement list just before the FOR_STMT
+     instead of in FOR_INIT_STMT.  If FOR_INIT_STMT happens to exist
+     but isn't a MODIFY_EXPR, give up.
+     handle it.  */
+
+  else if (init != NULL_TREE && TREE_CODE (init) != MODIFY_EXPR)
+    do_not_annotate_loop (&loop_info, as_invalid_initializer, NULL_TREE);
+
+  /* Examine the statement before the loop to see if it is a
+     valid initializer.  It must be either a MODIFY_EXPR or VAR_DECL,
+     possibly wrapped in language-specific structure.  */
+  else if (init == NULL_TREE && prev_tsi != NULL && tsi_stmt (*prev_tsi))
+    {
+      prev_stmt = tsi_stmt (*prev_tsi);
+
+      /* Call the language-specific hook to unwrap prev_stmt.  */
+      prev_stmt = (*lang_specific_unwrap_initializer) (prev_stmt);
+
+      /* See if we have a valid MODIFY_EXPR.  */
+      if (TREE_CODE (prev_stmt) == MODIFY_EXPR
+         && is_local_var (TREE_OPERAND (prev_stmt, 0))
+         && !TREE_SIDE_EFFECTS (TREE_OPERAND (prev_stmt, 1)))
+       {
+         decl = TREE_OPERAND (prev_stmt, 0);
+         init = prev_stmt;
+         unlink_prev = true;
+       }
+      else if (is_local_var (prev_stmt)
+              && !TREE_SIDE_EFFECTS (DECL_INITIAL (prev_stmt)))
+       {
+         /* If the preceding statement is the declaration of the loop
+            variable with its initialization, build an assignment
+            expression for the loop's initializer.  */
+         decl = prev_stmt;
+         init = build2 (MODIFY_EXPR, TREE_TYPE (decl), decl,
+                        DECL_INITIAL (decl));
+         /* We need to remove the initializer from the decl if we
+            end up using the init we just built instead.  */
+         fix_decl = true;
+       }
+    }
+
+  if (init == NULL_TREE || decl == NULL_TREE)
+    /* There is nothing we can do to find the correct init statement for
+       this loop.  */
+    do_not_annotate_loop (&loop_info, as_missing_initializer, NULL_TREE);
+
+  /* The condition must be a comparison of the decl we found in
+     the initializer against an expression that can be hoisted
+     outside the loop.  */
+  if (loop_info.state > as_in_kernels_loop)
+    /* Skip validating condition if we've already got an error.  */
+    ;
   else if (cond == NULL_TREE)
     do_not_annotate_loop (&loop_info, as_missing_predicate, NULL_TREE);
   else if (TREE_CODE_CLASS (TREE_CODE (cond)) != tcc_comparison)
     do_not_annotate_loop (&loop_info, as_invalid_predicate, cond);
   else
     {
-      /* The condition's LHS must be a local variable that does not
-        have its address taken.  Its RHS must also be such a local
-        variable or a constant.  */
-      tree induction_var = TREE_OPERAND (cond, 0);
-      tree limit_var = TREE_OPERAND (cond, 1);
-      if (!is_local_var (induction_var)
-         || (!is_local_var (limit_var)
-             && (TREE_CODE_CLASS (TREE_CODE (limit_var))
-                 != tcc_constant)))
+      tree limit_exp = NULL_TREE;
+
+      if (TREE_OPERAND (cond, 0) == decl)
+       limit_exp = TREE_OPERAND (cond, 1);
+      else if (TREE_OPERAND (cond, 1) == decl)
+       limit_exp = TREE_OPERAND (cond, 0);
+
+      if (!limit_exp
+         || (!is_local_var (limit_exp)
+             && (TREE_CODE_CLASS (TREE_CODE (limit_exp)) != tcc_constant)))
        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, induction_var,
-                                     loop_info.vars);
-         if (TREE_CODE_CLASS (TREE_CODE (limit_var)) != tcc_constant)
-           loop_info.vars = tree_cons (NULL_TREE, limit_var, loop_info.vars);
+         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);
        }
     }

@@ -3369,11 +3356,24 @@ check_and_annotate_for_loop (tree *nodeptr, tree_stmt_iterator *prev_tsi,
       /* If the traversal of the loop and all nested loops didn't hit
         any problems, attempt the actual transformation.  If it
         succeeds, replace this node with the annotated loop.  */
-      tree result = annotate_for_loop (node, prev_tsi, &loop_info);
+      tree result = annotate_for_loop (node, decl, init, &loop_info);
       if (result != node)
        {
          /* Success!  */
          *nodeptr = result;
+
+         if (unlink_prev)
+           /* We don't need the previous statement that we consumed
+              as an initializer in the new OMP_FOR any more.  */
+           tsi_delink (prev_tsi);
+
+         if (fix_decl)
+           /* We no longer need the initializer expression on the
+              decl of the loop variable and don't want to duplicate
+              it.  The kernels conversion pass would interpret it as
+              a stray assignment in a gang-single region.  */
+           DECL_INITIAL (decl) = NULL_TREE;
+
          return;
        }
     }
--
2.33.0

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