Hi all,
The attached patch fixes a bug in recent work on OpenACC "kernels"
regions. Jumps within nested binds or try statements were not analyzed
correctly and could lead to ICEs.
Tested on x86_64 with offloading to NVPTX.
Thanks,
Gergö
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.
>From 6db1f381f344b4482dcca6b82fc6316d172840be Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <ge...@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
