This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4, committed, PR65419] Add IFN_GOACC_DATA_END_WITH_ARG
- From: Tom de Vries <Tom_deVries at mentor dot com>
- To: <gcc-patches at gcc dot gnu dot org>
- Cc: Jakub Jelinek <jakub at redhat dot com>
- Date: Thu, 28 May 2015 17:27:47 +0200
- Subject: [gomp4, committed, PR65419] Add IFN_GOACC_DATA_END_WITH_ARG
- Authentication-results: sourceware.org; auth=none
Hi,
this patch fixes PR65419.
Consider this test-case:
...
void
f (void)
{
int i;
#pragma acc data copyout (i)
{
}
}
...
When compiling the oacc data region, the start and end are marked with
GOACC_data_start and GOACC_data_end:
...
.omp_data_arr.1.i = &i;
GOACC_data_start (-1, 1, &.omp_data_arr.1, &.omp_data_sizes.2,
&.omp_data_kinds.3);
GOACC_data_end ();
.omp_data_arr.1 = {CLOBBER};
..
We're marking the &.omp_data_arr.1 argument of GOACC_data_start with
fnspec 'r', meaning NOESCAPE and NOCLOBBER, which has the effect that
the call to GOACC_data_end is optimized to a tail call.
But actually, during GOACC_data_end we write i's accelerator value back
to i, which due to the tail call optimization is no longer allocated.
This causes a runtime error.
So actually, the fact that we write i's accelerator value back to i
during GOACC_data_end, means i and .omp_data_arr escape during
GOACC_data_start.
The easy way to fix this is to remove the 'r' in the fnspec for the
GOACC_data_start &.omp_data_arr. argument. But that would mean that
GOACC_data_start would become an optimization barrier, which would mean
missed optimizations in the kernels region.
This patch fixes the problem by adding the &.omp_data_arr argument to
the new internal function IFN_GOACC_DATA_END_WITH_ARG:
...
.omp_data_arr.1.i = &i;
GOACC_data_start (-1, 1, &.omp_data_arr.1, &.omp_data_sizes.2,
&.omp_data_kinds.3);
GOACC_DATA_END_WITH_ARG (&.omp_data_arr.1);
.omp_data_arr.1 = {CLOBBER};
...
This allows us to pretend that .omp_data_arr does not escape in
GOACC_data_start.
The internal function call is replaced by a GOACC_data_end call before
expand, dropping the argument not to break the abi:
...
.omp_data_arr.1.i = &i;
GOACC_data_start (-1, 1, &.omp_data_arr.1, &.omp_data_sizes.2,
&.omp_data_kinds.3);
GOACC_data_end ();
.omp_data_arr.1 ={v} {CLOBBER};
...
Bootstrapped and regtested on gomp-4_0-branch, committed to gomp-4_0-branch.
Thanks,
- Tom
Add IFN_GOACC_DATA_END_WITH_ARG
2015-05-28 Tom de Vries <tom@codesourcery.com>
PR tree-optimization/65419
* cfgexpand.c (pass_data_expand): Add PROP_gimple_lompifn to
properties_required field.
* gimplify.c (gimplify_omp_workshare): Use IFN_GOACC_DATA_END_WITH_ARG
instead of BUILT_IN_GOACC_DATA_END. Clear PROP_gimple_lompifn in
curr_properties.
(gimplify_function_tree): Tentatively set PROP_gimple_lompifn in
curr_properties.
* internal-fn.c (expand_GOACC_DATA_END_WITH_ARG): New dummy function.
* internal-fn.def (GOACC_DATA_END_WITH_ARG): New DEF_INTERNAL_FN.
* omp-low.c (lower_omp_target): Set argument of GOACC_DATA_END_WITH_ARG.
(pass_data_late_lower_omp): New pass_data.
(pass_late_lower_omp): New pass.
(pass_late_lower_omp::gate, pass_late_lower_omp::execute)
(make_pass_late_lower_omp): New function.
* passes.def: Add pass_late_lower_omp.
* tree-inline.c (expand_call_inline): Handle PROP_gimple_lompifn.
* tree-pass.h (PROP_gimple_lompifn): Add define.
* testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c: New test.
---
gcc/cfgexpand.c | 3 +-
gcc/gimplify.c | 25 +++++--
gcc/internal-fn.c | 9 +++
gcc/internal-fn.def | 1 +
gcc/omp-low.c | 86 +++++++++++++++++++++-
gcc/passes.def | 1 +
gcc/tree-inline.c | 16 ++--
gcc/tree-pass.h | 2 +
.../libgomp.oacc-c-c++-common/goacc-data-end.c | 68 +++++++++++++++++
9 files changed, 197 insertions(+), 14 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c
diff --git a/gcc/cfgexpand.c b/gcc/cfgexpand.c
index 5905ddb..6941e3e 100644
--- a/gcc/cfgexpand.c
+++ b/gcc/cfgexpand.c
@@ -5900,7 +5900,8 @@ const pass_data pass_data_expand =
( PROP_ssa | PROP_gimple_leh | PROP_cfg
| PROP_gimple_lcx
| PROP_gimple_lvec
- | PROP_gimple_lva), /* properties_required */
+ | PROP_gimple_lva
+ | PROP_gimple_lompifn), /* properties_required */
PROP_rtl, /* properties_provided */
( PROP_ssa | PROP_trees ), /* properties_destroyed */
0, /* todo_flags_start */
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 912b60f..c85b424 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7640,20 +7640,32 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (NULL);
if (ort == ORT_TARGET_DATA)
{
- enum built_in_function end_ix;
switch (TREE_CODE (expr))
{
case OACC_DATA:
- end_ix = BUILT_IN_GOACC_DATA_END;
+ /* Rather than building a call to BUILT_IN_GOACC_DATA_END, we use
+ this ifn which is similar, but has a pointer argument, which
+ will be later set to the &.omp_data_arr of the corresponding
+ BUILT_IN_GOACC_DATA_START.
+ This allows us to pretend that the &.omp_data_arr argument of
+ BUILT_IN_GOACC_DATA_START does not escape. */
+ g = gimple_build_call_internal (IFN_GOACC_DATA_END_WITH_ARG, 1,
+ null_pointer_node);
+ /* Clear the tentatively set PROP_gimple_lompifn, to indicate that
+ IFN_GOACC_DATA_END_WITH_ARG needs to be expanded. The argument
+ is not abi-compatible with the GOACC_data_end function, which
+ has no arguments. */
+ cfun->curr_properties &= ~PROP_gimple_lompifn;
break;
case OMP_TARGET_DATA:
- end_ix = BUILT_IN_GOMP_TARGET_END_DATA;
+ {
+ tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+ g = gimple_build_call (fn, 0);
+ }
break;
default:
gcc_unreachable ();
}
- tree fn = builtin_decl_explicit (end_ix);
- g = gimple_build_call (fn, 0);
gimple_seq cleanup = NULL;
gimple_seq_add_stmt (&cleanup, g);
g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
@@ -9484,6 +9496,9 @@ gimplify_function_tree (tree fndecl)
if necessary. */
cfun->curr_properties |= PROP_gimple_lva;
+ /* Tentatively set PROP_gimple_lompifn. */
+ cfun->curr_properties |= PROP_gimple_lompifn;
+
for (parm = DECL_ARGUMENTS (fndecl); parm ; parm = DECL_CHAIN (parm))
{
/* Preliminarily mark non-addressed complex variables as eligible
diff --git a/gcc/internal-fn.c b/gcc/internal-fn.c
index 0053ed9..27d05c7 100644
--- a/gcc/internal-fn.c
+++ b/gcc/internal-fn.c
@@ -1981,6 +1981,15 @@ expand_VA_ARG (gcall *stmt ATTRIBUTE_UNUSED)
gcc_unreachable ();
}
+/* GOACC_DATA_END_WITH_ARG is supposed to be expanded at pass_late_lower_omp.
+ So this dummy function should never be called. */
+
+static void
+expand_GOACC_DATA_END_WITH_ARG (gcall *stmt ATTRIBUTE_UNUSED)
+{
+ gcc_unreachable ();
+}
+
/* Routines to expand each internal function, indexed by function number.
Each routine has the prototype:
diff --git a/gcc/internal-fn.def b/gcc/internal-fn.def
index ba5c2c1..abe5c37 100644
--- a/gcc/internal-fn.def
+++ b/gcc/internal-fn.def
@@ -63,3 +63,4 @@ DEF_INTERNAL_FN (SUB_OVERFLOW, ECF_CONST | ECF_LEAF | ECF_NOTHROW, NULL)
DEF_INTERNAL_FN (MUL_OVERFLOW, ECF_CONST | ECF_LEAF | ECF_NOTHROW, NULL)
DEF_INTERNAL_FN (TSAN_FUNC_EXIT, ECF_NOVOPS | ECF_LEAF | ECF_NOTHROW, NULL)
DEF_INTERNAL_FN (VA_ARG, ECF_NOTHROW | ECF_LEAF, NULL)
+DEF_INTERNAL_FN (GOACC_DATA_END_WITH_ARG, ECF_NOTHROW, ".r")
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a9fd016..a3683a3 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12305,6 +12305,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
location_t loc = gimple_location (stmt);
bool offloaded, data_region, has_reduction;
unsigned int map_cnt = 0;
+ gimple goacc_data_end = NULL;
offloaded = is_gimple_omp_offloaded (stmt);
switch (gimple_omp_target_kind (stmt))
@@ -12336,7 +12337,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tgt_body = gimple_bind_body (tgt_bind);
}
else if (data_region)
- tgt_body = gimple_omp_body (stmt);
+ {
+ tgt_body = gimple_omp_body (stmt);
+ gimple try_stmt = gimple_seq_first_stmt (tgt_body);
+ gcc_assert (gimple_try_kind (try_stmt) == GIMPLE_TRY_FINALLY);
+ gimple_seq cleanup = gimple_try_cleanup (try_stmt);
+ if (gimple_call_internal_p (cleanup)
+ && gimple_call_internal_fn (cleanup) == IFN_GOACC_DATA_END_WITH_ARG)
+ {
+ goacc_data_end = cleanup;
+ gcc_assert (gimple_call_arg (goacc_data_end, 0) == null_pointer_node);
+ }
+ }
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
@@ -12469,6 +12481,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
= create_tmp_var (ctx->record_type, ".omp_data_arr");
DECL_NAMELESS (ctx->sender_decl) = 1;
TREE_ADDRESSABLE (ctx->sender_decl) = 1;
+
+ if (goacc_data_end != NULL)
+ {
+ tree arg = build_fold_addr_expr (ctx->sender_decl);
+ gimple_call_set_arg (goacc_data_end, 0, arg);
+ }
+
t = make_tree_vec (3);
TREE_VEC_ELT (t, 0) = ctx->sender_decl;
TREE_VEC_ELT (t, 1)
@@ -15012,4 +15031,69 @@ loop_in_oacc_kernels_region_p (struct loop *loop, basic_block *region_entry,
return false;
}
+namespace {
+
+const pass_data pass_data_late_lower_omp =
+{
+ GIMPLE_PASS, /* type */
+ "lateomplower", /* name */
+ OPTGROUP_NONE, /* optinfo_flags */
+ TV_NONE, /* tv_id */
+ ( PROP_cfg | PROP_ssa ), /* properties_required */
+ PROP_gimple_lompifn, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ 0, /* todo_flags_finish */
+};
+
+class pass_late_lower_omp : public gimple_opt_pass
+{
+public:
+ pass_late_lower_omp (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_late_lower_omp, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual unsigned int execute (function *);
+
+ virtual bool gate (function *)
+ {
+ return (cfun->curr_properties & PROP_gimple_lompifn) == 0;
+ }
+
+}; // class pass_lower_omp
+
+unsigned int
+pass_late_lower_omp::execute (function *fun)
+{
+ basic_block bb;
+ gimple_stmt_iterator i;
+
+ FOR_EACH_BB_FN (bb, fun)
+ for (i = gsi_start_bb (bb); !gsi_end_p (i); gsi_next (&i))
+ {
+ gimple stmt = gsi_stmt (i);
+ if (!(is_gimple_call (stmt)
+ && gimple_call_internal_p (stmt)
+ && gimple_call_internal_fn (stmt) == IFN_GOACC_DATA_END_WITH_ARG))
+ continue;
+
+ tree fn = builtin_decl_explicit (BUILT_IN_GOACC_DATA_END);
+ gimple g = gimple_build_call (fn, 0);
+
+ gsi_replace (&i, g, false);
+ }
+
+ return TODO_update_ssa;
+}
+
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_late_lower_omp (gcc::context *ctxt)
+{
+ return new pass_late_lower_omp (ctxt);
+}
+
#include "gt-omp-low.h"
diff --git a/gcc/passes.def b/gcc/passes.def
index 545287b..da497ed 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -358,6 +358,7 @@ along with GCC; see the file COPYING3. If not see
NEXT_PASS (pass_tm_edges);
POP_INSERT_PASSES ()
NEXT_PASS (pass_vtable_verify);
+ NEXT_PASS (pass_late_lower_omp);
NEXT_PASS (pass_lower_vaarg);
NEXT_PASS (pass_lower_vector);
NEXT_PASS (pass_lower_complex_O0);
diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 71d75d9..e6fef60 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -4525,13 +4525,15 @@ expand_call_inline (basic_block bb, gimple stmt, copy_body_data *id)
id->src_cfun = DECL_STRUCT_FUNCTION (fn);
id->call_stmt = stmt;
- /* If the the src function contains an IFN_VA_ARG, then so will the dst
- function after inlining. */
- if ((id->src_cfun->curr_properties & PROP_gimple_lva) == 0)
- {
- struct function *dst_cfun = DECL_STRUCT_FUNCTION (id->dst_fn);
- dst_cfun->curr_properties &= ~PROP_gimple_lva;
- }
+ {
+ /* Handle properties that need to be false in the resulting function, if
+ they're false in the src function. */
+ unsigned int props_mask = PROP_gimple_lva | PROP_gimple_lompifn;
+ unsigned int src_props = id->src_cfun->curr_properties;
+ unsigned int kill_props = props_mask & ~src_props;
+ struct function *dst_cfun = DECL_STRUCT_FUNCTION (id->dst_fn);
+ dst_cfun->curr_properties &= ~kill_props;
+ }
gcc_assert (!id->src_cfun->after_inlining);
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index 789dc64..6c79255 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -222,6 +222,7 @@ protected:
#define PROP_gimple_lvec (1 << 12) /* lowered vector */
#define PROP_gimple_eomp (1 << 13) /* no OpenMP directives */
#define PROP_gimple_lva (1 << 14) /* No va_arg internal function. */
+#define PROP_gimple_lompifn (1 << 15) /* No omp internal function. */
#define PROP_trees \
(PROP_gimple_any | PROP_gimple_lcf | PROP_gimple_leh | PROP_gimple_lomp)
@@ -403,6 +404,7 @@ extern gimple_opt_pass *make_pass_lower_complex (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_lower_vector (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_lower_vector_ssa (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_lower_omp (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_late_lower_omp (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_expand_omp_ssa (gcc::context *ctxt);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c
new file mode 100644
index 0000000..d3306aa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c
@@ -0,0 +1,68 @@
+/* { dg-do run } */
+/* { dg-options "-O2" } */
+
+/* Data directive at end of function. This is the variant that triggered
+ PR65419. */
+
+static void __attribute__((noinline,noclone))
+f (void)
+{
+ int i;
+
+#pragma acc data copyout (i)
+ {
+
+ }
+}
+
+/* Data directive in inlined function g_1. */
+
+static inline void
+g_1 (void)
+{
+ int i;
+
+#pragma acc data copyout (i)
+ {
+
+ }
+}
+
+static void __attribute__((noinline,noclone))
+g (void)
+{
+ g_1 ();
+}
+
+/* Data directive in function h into which a function h_1 is inlined. */
+
+static inline void
+h_1 (void)
+{
+
+}
+
+static void __attribute__((noinline,noclone))
+h (void)
+{
+ int i;
+
+ h_1 ();
+
+#pragma acc data copyout (i)
+ {
+
+ }
+}
+
+/* Main function calling the tests. */
+
+int
+main (void)
+{
+ f ();
+ g ();
+ h ();
+
+ return 0;
+}
--
1.9.1