Hi,
I've reverted the patch that added IFN_GOACC_DATA_END_WITH_ARG (
https://gcc.gnu.org/ml/gcc-patches/2015-05/msg02661.html ).
The patch attempted to fix a test failure, while at the same time
keeping the GOACC_data_start fnspec attributes to prevent it from
becoming an alias analysis optimization barrier.
Now that we've got -foffload-alias, we're no longer concerned about
GOACC builtins being alias analysis optimization barriers, so the
IFN_GOACC_DATA_END_WITH_ARG patch has become obsolete.
Committed to gomp-4_0-branch.
Thanks,
- Tom
Revert "Add IFN_GOACC_DATA_END_WITH_ARG"
2015-10-05 Tom de Vries <t...@codesourcery.com>
revert:
2015-05-28 Tom de Vries <t...@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 | 67 -----------------
9 files changed, 14 insertions(+), 196 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c
diff --git a/gcc/cfgexpand.c b/gcc/cfgexpand.c
index ca52d3d..bfbc958 100644
--- a/gcc/cfgexpand.c
+++ b/gcc/cfgexpand.c
@@ -6060,8 +6060,7 @@ const pass_data pass_data_expand =
( PROP_ssa | PROP_gimple_leh | PROP_cfg
| PROP_gimple_lcx
| PROP_gimple_lvec
- | PROP_gimple_lva
- | PROP_gimple_lompifn), /* properties_required */
+ | PROP_gimple_lva), /* 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 6283f0c..a5e28b4 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8960,32 +8960,20 @@ 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:
- /* 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;
+ end_ix = BUILT_IN_GOACC_DATA_END;
break;
case OMP_TARGET_DATA:
- {
- tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
- g = gimple_build_call (fn, 0);
- }
+ end_ix = BUILT_IN_GOMP_TARGET_END_DATA;
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);
@@ -10939,9 +10927,6 @@ 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 317149e..6fac752 100644
--- a/gcc/internal-fn.c
+++ b/gcc/internal-fn.c
@@ -1949,15 +1949,6 @@ 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 ();
-}
-
/* Expand the IFN_UNIQUE function according to its first argument. */
static void
diff --git a/gcc/internal-fn.def b/gcc/internal-fn.def
index 524b98d..ca06b10 100644
--- a/gcc/internal-fn.def
+++ b/gcc/internal-fn.def
@@ -65,7 +65,6 @@ 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")
/* An unduplicable, uncombinable function. Generally used to preserve
a CFG property in the face of jump threading, tail merging or
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index debedb1..130728a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -14999,7 +14999,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
bool offloaded, data_region;
unsigned int map_cnt = 0;
bool has_depend = false;
- gimple *goacc_data_end = NULL;
offloaded = is_gimple_omp_offloaded (stmt);
switch (gimple_omp_target_kind (stmt))
@@ -15044,18 +15043,7 @@ 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);
- 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);
- }
- }
+ tgt_body = gimple_omp_body (stmt);
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
@@ -15296,13 +15284,6 @@ 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)
@@ -18547,71 +18528,6 @@ omp_finish_file (void)
}
}
-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);
-}
-
/* Transform oacc_dim_size and oacc_dim_pos internal function calls to
constants, where possible. */
diff --git a/gcc/passes.def b/gcc/passes.def
index e68228e..5683bb7 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -367,7 +367,6 @@ along with GCC; see the file COPYING3. If not see
POP_INSERT_PASSES ()
NEXT_PASS (pass_simduid_cleanup);
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 00ad2af..3d06e6e 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -4540,15 +4540,13 @@ expand_call_inline (basic_block bb, gimple *stmt, copy_body_data *id)
id->src_cfun = DECL_STRUCT_FUNCTION (fn);
id->call_stmt = stmt;
- {
- /* 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;
- }
+ /* 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;
+ }
gcc_assert (!id->src_cfun->after_inlining);
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index 248c406..8eaf678 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -226,7 +226,6 @@ protected:
of math functions; the
current choices have
been optimized. */
-#define PROP_gimple_lompifn (1 << 16) /* No omp internal function. */
#define PROP_trees \
(PROP_gimple_any | PROP_gimple_lcf | PROP_gimple_leh | PROP_gimple_lomp)
@@ -415,7 +414,6 @@ 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
deleted file mode 100644
index 5685dbd..0000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/goacc-data-end.c
+++ /dev/null
@@ -1,67 +0,0 @@
-/* { dg-do run } */
-
-/* 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