[OpenACC,og8] OpenACC kernels control flow analysis bug fix

From 6db1f381f344b4482dcca6b82fc6316d172840be Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <gergo@codesourcery.com>
Date: Mon, 11 Feb 2019 08:23:31 -0800
Subject: [PATCH] OpenACC kernels control flow analysis bug fix

Correctly handle nested bind and try statements in the OpenACC kernels
conversion control-flow region analysis.

    gcc/
    * omp-oacc-kernels.c (control_flow_regions::compute_regions): Factored
    out...
    (control_flow_regions::visit_gimple_seq): ... this new method, now also
    handling bind and try statements.

    gcc/testsuite/
    * gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c: Add tests.
---
 gcc/omp-oacc-kernels.c                             | 94 +++++++++++++++-------
 .../c-c++-common/goacc/kernels-decompose-1.c       | 44 ++++++++++
 2 files changed, 107 insertions(+), 31 deletions(-)

diff --git a/gcc/omp-oacc-kernels.c b/gcc/omp-oacc-kernels.c
index d1db492..1fa2647 100644
--- a/gcc/omp-oacc-kernels.c
+++ b/gcc/omp-oacc-kernels.c
@@ -935,12 +935,24 @@  class control_flow_regions
        control-flow regions in the statement sequence SEQ.  */
     void compute_regions (gimple_seq seq);
 
+    /* Helper for compute_regions, scanning a single statement sequence SEQ
+       starting at index IDX and returning the next index after the last
+       statement in the sequence.  */
+    size_t visit_gimple_seq (gimple_seq seq, size_t idx);
+
     /* The mapping from statement indices to region representatives.  */
     vec <size_t> representatives;
 
     /* A cache mapping statement indices to a flag indicating whether the
        statement is a top level OpenACC for loop.  */
     vec <bool> omp_for_loops;
+
+    /* A mapping of control flow statements (goto, switch, cond) to their
+       representatives.  */
+    hash_map <gimple *, size_t> control_flow_reps;
+
+    /* A mapping of labels to their representatives.  */
+    hash_map <tree, size_t> label_reps;
 };
 
 control_flow_regions::control_flow_regions (gimple_seq seq)
@@ -1008,41 +1020,12 @@  control_flow_regions::union_reps (size_t a, size_t b)
 void
 control_flow_regions::compute_regions (gimple_seq seq)
 {
-  hash_map <gimple *, size_t> control_flow_reps;
-  hash_map <tree, size_t> label_reps;
-  size_t current_region = 0, idx = 0;
+  size_t idx = 0;
 
   /* In a first pass, assign an initial region to each statement.  Except in
      the case of OpenACC loops, each statement simply gets the same region
      representative as its predecessor.  */
-  for (gimple_stmt_iterator gsi = gsi_start (seq);
-       !gsi_end_p (gsi);
-       gsi_next (&gsi))
-    {
-      gimple *stmt = gsi_stmt (gsi);
-      gimple *omp_for = top_level_omp_for_in_stmt (stmt);
-      omp_for_loops.safe_push (omp_for != NULL);
-      if (omp_for != NULL)
-        {
-          /* Assign a new region to this loop and to its successor.  */
-          current_region = idx;
-          representatives.safe_push (current_region);
-          current_region++;
-        }
-      else
-        {
-          representatives.safe_push (current_region);
-          /* Remember any jumps and labels for the second pass below.  */
-          if (gimple_code (stmt) == GIMPLE_COND
-              || gimple_code (stmt) == GIMPLE_SWITCH
-              || gimple_code (stmt) == GIMPLE_GOTO)
-            control_flow_reps.put (stmt, current_region);
-          else if (gimple_code (stmt) == GIMPLE_LABEL)
-            label_reps.put (gimple_label_label (as_a <glabel *> (stmt)),
-                            current_region);
-        }
-      idx++;
-    }
+  visit_gimple_seq (seq, idx);
   gcc_assert (representatives.length () == omp_for_loops.length ());
 
   /* Revisit all the control flow statements and union the region of each
@@ -1087,6 +1070,55 @@  control_flow_regions::compute_regions (gimple_seq seq)
     }
 }
 
+size_t
+control_flow_regions::visit_gimple_seq (gimple_seq seq, size_t idx)
+{
+  size_t current_region = idx;
+
+  for (gimple_stmt_iterator gsi = gsi_start (seq);
+       !gsi_end_p (gsi);
+       gsi_next (&gsi))
+    {
+      gimple *stmt = gsi_stmt (gsi);
+      gimple *omp_for = top_level_omp_for_in_stmt (stmt);
+      omp_for_loops.safe_push (omp_for != NULL);
+      if (omp_for != NULL)
+        {
+          /* Assign a new region to this loop and to its successor.  */
+          current_region = idx;
+          representatives.safe_push (current_region);
+          current_region++;
+        }
+      else if (gimple_code (stmt) == GIMPLE_BIND)
+        {
+          representatives.safe_push (current_region);
+          gimple_seq body = gimple_bind_body (as_a <gbind *> (stmt));
+          idx = visit_gimple_seq (body, idx);
+        }
+      else if (gimple_code (stmt) == GIMPLE_TRY)
+        {
+          representatives.safe_push (current_region);
+          idx = visit_gimple_seq (gimple_try_eval (stmt), idx);
+          idx = visit_gimple_seq (gimple_try_cleanup (stmt), idx);
+        }
+      else
+        {
+          representatives.safe_push (current_region);
+          /* Remember any jumps and labels for the second pass below.  */
+          if (gimple_code (stmt) == GIMPLE_COND
+              || gimple_code (stmt) == GIMPLE_SWITCH
+              || gimple_code (stmt) == GIMPLE_GOTO)
+            control_flow_reps.put (stmt, current_region);
+          else if (gimple_code (stmt) == GIMPLE_LABEL)
+            label_reps.put (gimple_label_label (as_a <glabel *> (stmt)),
+                            current_region);
+        }
+      idx++;
+    }
+
+  return idx;
+}
+
 /* Decompose the body of the KERNELS_REGION, which was originally annotated
    with the KERNELS_CLAUSES, into a series of regions.  */
 
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
index 7255d69..6a3381e 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
@@ -118,5 +118,49 @@  main ()
 #pragma acc kernels /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */
   ;
 
+  #pragma acc kernels
+  {
+    int a[20], b[20];
+    int n = 0;
+
+    /* The goto in this loop is placed into a bind; make sure it is
+       traversed correctly.  */
+    for (int i = 0; i < 20; i++)
+      goto label_43;
+
+    #pragma acc loop independent /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" } */
+    for (int m = 0; m < 20; m++)
+      b[m] = a[m];
+
+label_43:
+    b[n] = a[n];
+  }
+
+#ifdef __cplusplus
+  #pragma acc kernels
+  {
+    int a[20], b[20];
+    int n = 0;
+
+    for (int i = 0; i < 20; i++)
+      {
+        try
+          {
+            goto label_44;
+          }
+        catch (int)
+          {
+          }
+      }
+
+    #pragma acc loop independent /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target c++ } } */
+    for (int m = 0; m < 20; m++)
+      b[m] = a[m];
+
+label_44:
+    b[n] = a[n];
+  }
+#endif
+
   return 0;
 }
-- 
2.8.1