-/* Lowering pass for OpenMP directives. Converts OpenMP directives
- into explicit calls to the runtime library (libgomp) and data
- marshalling to implement data sharing and copying clauses.
+/* Lowering pass for OMP directives. Converts OMP directives into explicit
+ calls to the runtime library (libgomp), data marshalling to implement data
+ sharing and copying clauses, offloading to accelerators, and more.
+
Contributed by Diego Novillo <dnovillo@redhat.com>
- Copyright (C) 2005-2014 Free Software Foundation, Inc.
+ Copyright (C) 2005-2015 Free Software Foundation, Inc.
This file is part of GCC.
#include "system.h"
#include "coretypes.h"
#include "tm.h"
+#include "alias.h"
+#include "symtab.h"
#include "tree.h"
+#include "fold-const.h"
#include "stringpool.h"
#include "stor-layout.h"
#include "rtl.h"
-#include "pointer-set.h"
+#include "predict.h"
+#include "hard-reg-set.h"
+#include "function.h"
+#include "dominance.h"
+#include "cfg.h"
+#include "cfganal.h"
#include "basic-block.h"
#include "tree-ssa-alias.h"
#include "internal-fn.h"
#include "gimple-fold.h"
#include "gimple-expr.h"
-#include "is-a.h"
#include "gimple.h"
#include "gimplify.h"
#include "gimple-iterator.h"
#include "ssa-iterators.h"
#include "tree-ssanames.h"
#include "tree-into-ssa.h"
+#include "flags.h"
+#include "insn-config.h"
+#include "expmed.h"
+#include "dojump.h"
+#include "explow.h"
+#include "calls.h"
+#include "emit-rtl.h"
+#include "varasm.h"
+#include "stmt.h"
#include "expr.h"
#include "tree-dfa.h"
#include "tree-ssa.h"
-#include "flags.h"
-#include "function.h"
-#include "expr.h"
#include "tree-pass.h"
#include "except.h"
#include "splay-tree.h"
+#include "insn-codes.h"
#include "optabs.h"
#include "cfgloop.h"
#include "target.h"
+#include "common/common-target.h"
#include "omp-low.h"
#include "gimple-low.h"
#include "tree-cfgcleanup.h"
#include "pretty-print.h"
+#include "alloc-pool.h"
+#include "symbol-summary.h"
#include "ipa-prop.h"
#include "tree-nested.h"
#include "tree-eh.h"
+#include "cilk.h"
+#include "context.h"
+#include "lto-section-names.h"
+#include "gomp-constants.h"
-/* Lowering of OpenMP parallel and workshare constructs proceeds in two
+/* Lowering of OMP parallel and workshare constructs proceeds in two
phases. The first phase scans the function looking for OMP statements
and then for variables that must be replaced to satisfy data sharing
clauses. The second phase expands code for the constructs, as well as
expressions.
Final code generation is done by pass_expand_omp. The flowgraph is
- scanned for parallel regions which are then moved to a new
- function, to be invoked by the thread library. */
+ scanned for regions which are then moved to a new
+ function, to be invoked by the thread library, or offloaded. */
-/* Parallel region information. Every parallel and workshare
+/* OMP region information. Every parallel and workshare
directive is enclosed between two markers, the OMP_* directive
and a corresponding OMP_RETURN statement. */
bool is_combined_parallel;
};
+/* Levels of parallelism as defined by OpenACC. Increasing numbers
+ correspond to deeper loop nesting levels. */
+#define MASK_GANG 1
+#define MASK_WORKER 2
+#define MASK_VECTOR 4
+
/* Context structure. Used to store information about each parallel
directive in the code. */
construct. In the case of a parallel, this is in the child function. */
tree block_vars;
+ /* A map of reduction pointer variables. For accelerators, each
+ reduction variable is replaced with an array. Each thread, in turn,
+ is assigned to a slot on that array. */
+ splay_tree reduction_map;
+
/* Label to which GOMP_cancel{,llation_point} and explicit and implicit
barriers should jump to during omplower pass. */
tree cancel_label;
/* True if this construct can be cancelled. */
bool cancellable;
+
+ /* For OpenACC loops, a mask of gang, worker and vector used at
+ levels below this one. */
+ int gwv_below;
+ /* For OpenACC loops, a mask of gang, worker and vector used at
+ this level and above. For parallel and kernels clauses, a mask
+ indicating which of num_gangs/num_workers/num_vectors was used. */
+ int gwv_this;
} omp_context;
+/* A structure holding the elements of:
+ for (V = N1; V cond N2; V += STEP) [...] */
struct omp_for_data_loop
{
{
struct omp_for_data_loop loop;
tree chunk_size;
- gimple for_stmt;
+ gomp_for *for_stmt;
tree pre, iter_type;
int collapse;
bool have_nowait, have_ordered;
static int target_nesting_level;
static struct omp_region *root_omp_region;
static bitmap task_shared_vars;
+static vec<omp_context *> taskreg_contexts;
static void scan_omp (gimple_seq *, omp_context *);
static tree scan_omp_1_op (tree *, int *, void *);
*handled_ops_p = false; \
break;
+/* Helper function to get the name of the array containing the partial
+ reductions for OpenACC reductions. */
+static const char *
+oacc_get_reduction_array_id (tree node)
+{
+ const char *id = IDENTIFIER_POINTER (DECL_NAME (node));
+ int len = strlen ("OACC") + strlen (id);
+ char *temp_name = XALLOCAVEC (char, len + 1);
+ snprintf (temp_name, len + 1, "OACC%s", id);
+ return IDENTIFIER_POINTER (get_identifier (temp_name));
+}
+
+/* Determine the number of threads OpenACC threads used to determine the
+ size of the array of partial reductions. Currently, this is num_gangs
+ * vector_length. This value may be different than GOACC_GET_NUM_THREADS,
+ because it is independed of the device used. */
+
+static tree
+oacc_max_threads (omp_context *ctx)
+{
+ tree nthreads, vector_length, gangs, clauses;
+
+ gangs = fold_convert (sizetype, integer_one_node);
+ vector_length = gangs;
+
+ /* The reduction clause may be nested inside a loop directive.
+ Scan for the innermost vector_length clause. */
+ for (omp_context *oc = ctx; oc; oc = oc->outer)
+ {
+ if (gimple_code (oc->stmt) != GIMPLE_OMP_TARGET
+ || (gimple_omp_target_kind (oc->stmt)
+ != GF_OMP_TARGET_KIND_OACC_PARALLEL))
+ continue;
+
+ clauses = gimple_omp_target_clauses (oc->stmt);
+
+ vector_length = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH);
+ if (vector_length)
+ vector_length = fold_convert_loc (OMP_CLAUSE_LOCATION (vector_length),
+ sizetype,
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR
+ (vector_length));
+ else
+ vector_length = fold_convert (sizetype, integer_one_node);
+
+ gangs = find_omp_clause (clauses, OMP_CLAUSE_NUM_GANGS);
+ if (gangs)
+ gangs = fold_convert_loc (OMP_CLAUSE_LOCATION (gangs), sizetype,
+ OMP_CLAUSE_NUM_GANGS_EXPR (gangs));
+ else
+ gangs = fold_convert (sizetype, integer_one_node);
+
+ break;
+ }
+
+ nthreads = fold_build2 (MULT_EXPR, sizetype, gangs, vector_length);
+
+ return nthreads;
+}
+
+/* Holds offload tables with decls. */
+vec<tree, va_gc> *offload_funcs, *offload_vars;
+
/* Convenience function for calling scan_omp_1_op on tree operands. */
static inline tree
static tree lookup_decl_in_outer_ctx (tree, omp_context *);
static tree maybe_lookup_decl_in_outer_ctx (tree, omp_context *);
-/* Find an OpenMP clause of type KIND within CLAUSES. */
+/* Find an OMP clause of type KIND within CLAUSES. */
tree
find_omp_clause (tree clauses, enum omp_clause_code kind)
them into *FD. */
static void
-extract_omp_for_data (gimple for_stmt, struct omp_for_data *fd,
+extract_omp_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
struct omp_for_data_loop *loops)
{
tree t, var, *collapse_iter, *collapse_count;
int i;
struct omp_for_data_loop dummy_loop;
location_t loc = gimple_location (for_stmt);
- bool simd = gimple_omp_for_kind (for_stmt) & GF_OMP_FOR_KIND_SIMD;
+ bool simd = gimple_omp_for_kind (for_stmt) & GF_OMP_FOR_SIMD;
bool distribute = gimple_omp_for_kind (for_stmt)
== GF_OMP_FOR_KIND_DISTRIBUTE;
fd->have_ordered = false;
fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
fd->chunk_size = NULL_TREE;
+ if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_CILKFOR)
+ fd->sched_kind = OMP_CLAUSE_SCHEDULE_CILKFOR;
collapse_iter = NULL;
collapse_count = NULL;
collapse_iter = &OMP_CLAUSE_COLLAPSE_ITERVAR (t);
collapse_count = &OMP_CLAUSE_COLLAPSE_COUNT (t);
}
+ break;
default:
break;
}
break;
case NE_EXPR:
gcc_assert (gimple_omp_for_kind (for_stmt)
- == GF_OMP_FOR_KIND_CILKSIMD);
+ == GF_OMP_FOR_KIND_CILKSIMD
+ || (gimple_omp_for_kind (for_stmt)
+ == GF_OMP_FOR_KIND_CILKFOR));
break;
case LE_EXPR:
if (POINTER_TYPE_P (TREE_TYPE (loop->n2)))
fd->loop.step = build_int_cst (TREE_TYPE (fd->loop.v), 1);
fd->loop.cond_code = LT_EXPR;
}
+
+ /* For OpenACC loops, force a chunk size of one, as this avoids the default
+ scheduling where several subsequent iterations are being executed by the
+ same thread. */
+ if (gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ {
+ gcc_assert (fd->chunk_size == NULL_TREE);
+ fd->chunk_size = build_int_cst (TREE_TYPE (fd->loop.v), 1);
+ }
}
gcc_assert (gimple_code (ws_stmt) == GIMPLE_OMP_FOR);
- extract_omp_for_data (ws_stmt, &fd, NULL);
+ extract_omp_for_data (as_a <gomp_for *> (ws_stmt), &fd, NULL);
if (fd.collapse > 1 && TREE_CODE (fd.loop.n2) != INTEGER_CST)
return false;
location_t loc = gimple_location (ws_stmt);
vec<tree, va_gc> *ws_args;
- if (gimple_code (ws_stmt) == GIMPLE_OMP_FOR)
+ if (gomp_for *for_stmt = dyn_cast <gomp_for *> (ws_stmt))
{
struct omp_for_data fd;
tree n1, n2;
- extract_omp_for_data (ws_stmt, &fd, NULL);
+ extract_omp_for_data (for_stmt, &fd, NULL);
n1 = fd.loop.n1;
n2 = fd.loop.n2;
- if (gimple_omp_for_combined_into_p (ws_stmt))
+ if (gimple_omp_for_combined_into_p (for_stmt))
{
tree innerc
= find_omp_clause (gimple_omp_parallel_clauses (par_stmt),
return lang_hooks.decls.omp_privatize_by_reference (decl);
}
-/* Lookup variables in the decl or field splay trees. The "maybe" form
+/* Return the type of a decl. If the decl is reference type,
+ return its base type. */
+static inline tree
+get_base_type (tree decl)
+{
+ tree type = TREE_TYPE (decl);
+ if (is_reference (decl))
+ type = TREE_TYPE (type);
+ return type;
+}
+
+/* Lookup variables. The "maybe" form
allows for the variable form to not have been entered, otherwise we
assert that the variable must have been entered. */
static inline tree
lookup_decl (tree var, omp_context *ctx)
{
- tree *n;
- n = (tree *) pointer_map_contains (ctx->cb.decl_map, var);
+ tree *n = ctx->cb.decl_map->get (var);
return *n;
}
static inline tree
maybe_lookup_decl (const_tree var, omp_context *ctx)
{
- tree *n;
- n = (tree *) pointer_map_contains (ctx->cb.decl_map, var);
+ tree *n = ctx->cb.decl_map->get (const_cast<tree> (var));
return n ? *n : NULL_TREE;
}
return n ? (tree) n->value : NULL_TREE;
}
+static inline tree
+lookup_oacc_reduction (const char *id, omp_context *ctx)
+{
+ splay_tree_node n;
+ n = splay_tree_lookup (ctx->reduction_map, (splay_tree_key) id);
+ return (tree) n->value;
+}
+
+static inline tree
+maybe_lookup_oacc_reduction (tree var, omp_context *ctx)
+{
+ splay_tree_node n = NULL;
+ if (ctx->reduction_map)
+ n = splay_tree_lookup (ctx->reduction_map, (splay_tree_key) var);
+ return n ? (tree) n->value : NULL_TREE;
+}
+
/* Return true if DECL should be copied by pointer. SHARED_CTX is
the parallel context if DECL is to be shared. */
when we know the value is not accessible from an outer scope. */
if (shared_ctx)
{
+ gcc_assert (!is_gimple_omp_oacc (shared_ctx->stmt));
+
/* ??? Trivially accessible from anywhere. But why would we even
be passing an address in this case? Should we simply assert
this to be false, or should we have a cleanup pass that removes
x = build_receiver_ref (var, by_ref, ctx);
}
else if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
{
/* #pragma omp simd isn't a worksharing construct, and can reference even
private vars in its linear etc. clauses. */
|| !splay_tree_lookup (ctx->field_map, (splay_tree_key) var));
gcc_assert ((mask & 2) == 0 || !ctx->sfield_map
|| !splay_tree_lookup (ctx->sfield_map, (splay_tree_key) var));
+ gcc_assert ((mask & 3) == 3
+ || !is_gimple_omp_oacc (ctx->stmt));
type = TREE_TYPE (var);
if (mask & 4)
ctx->cb = outer_ctx->cb;
ctx->cb.block = NULL;
ctx->depth = outer_ctx->depth + 1;
+ ctx->reduction_map = outer_ctx->reduction_map;
}
else
{
ctx->cb.src_fn = current_function_decl;
ctx->cb.dst_fn = current_function_decl;
- ctx->cb.src_node = cgraph_get_node (current_function_decl);
+ ctx->cb.src_node = cgraph_node::get (current_function_decl);
gcc_checking_assert (ctx->cb.src_node);
ctx->cb.dst_node = ctx->cb.src_node;
ctx->cb.src_cfun = cfun;
ctx->depth = 1;
}
- ctx->cb.decl_map = pointer_map_create ();
+ ctx->cb.decl_map = new hash_map<tree, tree>;
return ctx;
}
/* Finalize task copyfn. */
static void
-finalize_task_copyfn (gimple task_stmt)
+finalize_task_copyfn (gomp_task *task_stmt)
{
struct function *child_cfun;
tree child_fn;
gimple_seq seq = NULL, new_seq;
- gimple bind;
+ gbind *bind;
child_fn = gimple_omp_task_copy_fn (task_stmt);
if (child_fn == NULL_TREE)
pop_cfun ();
/* Inform the callgraph about the new function. */
- cgraph_add_new_function (child_fn, false);
+ cgraph_node *node = cgraph_node::get_create (child_fn);
+ node->parallelized_function = 1;
+ cgraph_node::add_new_function (child_fn, false);
}
/* Destroy a omp_context data structures. Called through the splay tree
{
omp_context *ctx = (omp_context *) value;
- pointer_map_destroy (ctx->cb.decl_map);
+ delete ctx->cb.decl_map;
if (ctx->field_map)
splay_tree_delete (ctx->field_map);
if (ctx->sfield_map)
splay_tree_delete (ctx->sfield_map);
+ /* Reduction map is copied to nested contexts, so only delete it in the
+ owner. */
+ if (ctx->reduction_map
+ && gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+ && is_gimple_omp_offloaded (ctx->stmt)
+ && is_gimple_omp_oacc (ctx->stmt))
+ splay_tree_delete (ctx->reduction_map);
/* We hijacked DECL_ABSTRACT_ORIGIN earlier. We need to clear it before
it produces corrupt debug information. */
}
if (is_task_ctx (ctx))
- finalize_task_copyfn (ctx->stmt);
+ finalize_task_copyfn (as_a <gomp_task *> (ctx->stmt));
XDELETE (ctx);
}
layout_type (type);
}
- TREE_TYPE (ctx->receiver_decl) = build_pointer_type (type);
+ TREE_TYPE (ctx->receiver_decl)
+ = build_qualified_type (build_reference_type (type), TYPE_QUAL_RESTRICT);
}
/* Instantiate decls as necessary in CTX to satisfy the data sharing
break;
case OMP_CLAUSE_SHARED:
+ decl = OMP_CLAUSE_DECL (c);
/* Ignore shared directives in teams construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
- break;
+ {
+ /* Global variables don't need to be copied,
+ the receiver side will use them directly. */
+ tree odecl = maybe_lookup_decl_in_outer_ctx (decl, ctx);
+ if (is_global_var (odecl))
+ break;
+ insert_decl_map (&ctx->cb, decl, odecl);
+ break;
+ }
gcc_assert (is_taskreg_ctx (ctx));
- decl = OMP_CLAUSE_DECL (c);
gcc_assert (!COMPLETE_TYPE_P (TREE_TYPE (decl))
|| !is_variable_sized (decl));
/* Global variables don't need to be copied,
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
+ /* FALLTHRU */
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
install_var_field (decl, by_ref, 3, ctx);
}
install_var_local (decl, ctx);
+ if (is_gimple_omp_oacc (ctx->stmt)
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ /* Create a decl for the reduction array. */
+ tree var = OMP_CLAUSE_DECL (c);
+ tree type = get_base_type (var);
+ tree ptype = build_pointer_type (type);
+ tree array = create_tmp_var (ptype,
+ oacc_get_reduction_array_id (var));
+ omp_context *c = (ctx->field_map ? ctx : ctx->outer);
+ install_var_field (array, true, 3, c);
+ install_var_local (array, c);
+
+ /* Insert it into the current context. */
+ splay_tree_insert (ctx->reduction_map, (splay_tree_key)
+ oacc_get_reduction_array_id (var),
+ (splay_tree_value) array);
+ splay_tree_insert (ctx->reduction_map,
+ (splay_tree_key) array,
+ (splay_tree_value) array);
+ }
break;
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE__CILK_FOR_COUNT_:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& DECL_P (decl)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
- && lookup_attribute ("omp declare target",
- DECL_ATTRIBUTES (decl)))
+ && varpool_node::get_create (decl)->offloadable)
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
{
- /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
- #pragma omp target data, there is nothing to map for
- those. */
- if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
- && !POINTER_TYPE_P (TREE_TYPE (decl)))
+ /* Ignore GOMP_MAP_POINTER kind for arrays in regions that are
+ not offloaded; there is nothing to map for those. */
+ if (!is_gimple_omp_offloaded (ctx->stmt)
+ && !POINTER_TYPE_P (TREE_TYPE (decl))
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
break;
}
if (DECL_P (decl))
else
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 7, ctx);
else
install_var_field (decl, true, 3, ctx);
- if (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_REGION)
+ if (is_gimple_omp_offloaded (ctx->stmt))
install_var_local (decl, ctx);
}
}
&& nc != NULL_TREE
&& OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_DECL (nc) == base
- && OMP_CLAUSE_MAP_KIND (nc) == OMP_CLAUSE_MAP_POINTER
+ && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER
&& integer_zerop (OMP_CLAUSE_SIZE (nc)))
{
OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) = 1;
}
else
{
+ if (ctx->outer)
+ {
+ scan_omp_op (&OMP_CLAUSE_DECL (c), ctx->outer);
+ decl = OMP_CLAUSE_DECL (c);
+ }
gcc_assert (!splay_tree_lookup (ctx->field_map,
(splay_tree_key) decl));
tree field
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
break;
case OMP_CLAUSE_ALIGNED:
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE__CACHE_:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
+ sorry ("Clause not supported yet");
+ break;
+
default:
gcc_unreachable ();
}
break;
/* FALLTHRU */
- case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
+ /* FALLTHRU */
+ case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
break;
case OMP_CLAUSE_MAP:
- if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+ if (!is_gimple_omp_offloaded (ctx->stmt))
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
- && lookup_attribute ("omp declare target",
- DECL_ATTRIBUTES (decl)))
+ && varpool_node::get_create (decl)->offloadable)
break;
if (DECL_P (decl))
{
- if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
&& !COMPLETE_TYPE_P (TREE_TYPE (decl)))
{
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE__CILK_FOR_COUNT_:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WAIT:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ break;
+
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE__CACHE_:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_SEQ:
+ sorry ("Clause not supported yet");
break;
default:
}
}
+ gcc_checking_assert (!scan_array_reductions
+ || !is_gimple_omp_oacc (ctx->stmt));
if (scan_array_reductions)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
scan_omp (&OMP_CLAUSE_LINEAR_GIMPLE_SEQ (c), ctx);
}
-/* Create a new name for omp child function. Returns an identifier. */
+/* Create a new name for omp child function. Returns an identifier. If
+ IS_CILK_FOR is true then the suffix for the child function is
+ "_cilk_for_fn." */
+
+static tree
+create_omp_child_function_name (bool task_copy, bool is_cilk_for)
+{
+ if (is_cilk_for)
+ return clone_function_name (current_function_decl, "_cilk_for_fn");
+ return clone_function_name (current_function_decl,
+ task_copy ? "_omp_cpyfn" : "_omp_fn");
+}
+
+/* Returns the type of the induction variable for the child function for
+ _Cilk_for and the types for _high and _low variables based on TYPE. */
static tree
-create_omp_child_function_name (bool task_copy)
+cilk_for_check_loop_diff_type (tree type)
{
- return (clone_function_name (current_function_decl,
- task_copy ? "_omp_cpyfn" : "_omp_fn"));
+ if (TYPE_PRECISION (type) <= TYPE_PRECISION (uint32_type_node))
+ {
+ if (TYPE_UNSIGNED (type))
+ return uint32_type_node;
+ else
+ return integer_type_node;
+ }
+ else
+ {
+ if (TYPE_UNSIGNED (type))
+ return uint64_type_node;
+ else
+ return long_long_integer_type_node;
+ }
}
/* Build a decl for the omp child function. It'll not contain a body
{
tree decl, type, name, t;
- name = create_omp_child_function_name (task_copy);
+ tree cilk_for_count
+ = (flag_cilkplus && gimple_code (ctx->stmt) == GIMPLE_OMP_PARALLEL)
+ ? find_omp_clause (gimple_omp_parallel_clauses (ctx->stmt),
+ OMP_CLAUSE__CILK_FOR_COUNT_) : NULL_TREE;
+ tree cilk_var_type = NULL_TREE;
+
+ name = create_omp_child_function_name (task_copy,
+ cilk_for_count != NULL_TREE);
if (task_copy)
type = build_function_type_list (void_type_node, ptr_type_node,
ptr_type_node, NULL_TREE);
+ else if (cilk_for_count)
+ {
+ type = TREE_TYPE (OMP_CLAUSE_OPERAND (cilk_for_count, 0));
+ cilk_var_type = cilk_for_check_loop_diff_type (type);
+ type = build_function_type_list (void_type_node, ptr_type_node,
+ cilk_var_type, cilk_var_type, NULL_TREE);
+ }
else
type = build_function_type_list (void_type_node, ptr_type_node, NULL_TREE);
- decl = build_decl (gimple_location (ctx->stmt),
- FUNCTION_DECL, name, type);
+ decl = build_decl (gimple_location (ctx->stmt), FUNCTION_DECL, name, type);
+ gcc_checking_assert (!is_gimple_omp_oacc (ctx->stmt)
+ || !task_copy);
if (!task_copy)
ctx->cb.dst_fn = decl;
else
TREE_STATIC (decl) = 1;
TREE_USED (decl) = 1;
DECL_ARTIFICIAL (decl) = 1;
- DECL_NAMELESS (decl) = 1;
DECL_IGNORED_P (decl) = 0;
TREE_PUBLIC (decl) = 0;
DECL_UNINLINABLE (decl) = 1;
DECL_EXTERNAL (decl) = 0;
DECL_CONTEXT (decl) = NULL_TREE;
DECL_INITIAL (decl) = make_node (BLOCK);
- bool target_p = false;
- if (lookup_attribute ("omp declare target",
- DECL_ATTRIBUTES (current_function_decl)))
- target_p = true;
+ if (cgraph_node::get (current_function_decl)->offloadable)
+ cgraph_node::get_create (decl)->offloadable = 1;
else
{
omp_context *octx;
for (octx = ctx; octx; octx = octx->outer)
- if (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
- && gimple_omp_target_kind (octx->stmt)
- == GF_OMP_TARGET_KIND_REGION)
+ if (is_gimple_omp_offloaded (octx->stmt))
{
- target_p = true;
+ cgraph_node::get_create (decl)->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+#endif
break;
}
}
- if (target_p)
+
+ if (cgraph_node::get_create (decl)->offloadable
+ && !lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (current_function_decl)))
DECL_ATTRIBUTES (decl)
- = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, DECL_ATTRIBUTES (decl));
+ = tree_cons (get_identifier ("omp target entrypoint"),
+ NULL_TREE, DECL_ATTRIBUTES (decl));
t = build_decl (DECL_SOURCE_LOCATION (decl),
RESULT_DECL, NULL_TREE, void_type_node);
DECL_CONTEXT (t) = decl;
DECL_RESULT (decl) = t;
- t = build_decl (DECL_SOURCE_LOCATION (decl),
- PARM_DECL, get_identifier (".omp_data_i"), ptr_type_node);
+ /* _Cilk_for's child function requires two extra parameters called
+ __low and __high that are set the by Cilk runtime when it calls this
+ function. */
+ if (cilk_for_count)
+ {
+ t = build_decl (DECL_SOURCE_LOCATION (decl),
+ PARM_DECL, get_identifier ("__high"), cilk_var_type);
+ DECL_ARTIFICIAL (t) = 1;
+ DECL_NAMELESS (t) = 1;
+ DECL_ARG_TYPE (t) = ptr_type_node;
+ DECL_CONTEXT (t) = current_function_decl;
+ TREE_USED (t) = 1;
+ DECL_CHAIN (t) = DECL_ARGUMENTS (decl);
+ DECL_ARGUMENTS (decl) = t;
+
+ t = build_decl (DECL_SOURCE_LOCATION (decl),
+ PARM_DECL, get_identifier ("__low"), cilk_var_type);
+ DECL_ARTIFICIAL (t) = 1;
+ DECL_NAMELESS (t) = 1;
+ DECL_ARG_TYPE (t) = ptr_type_node;
+ DECL_CONTEXT (t) = current_function_decl;
+ TREE_USED (t) = 1;
+ DECL_CHAIN (t) = DECL_ARGUMENTS (decl);
+ DECL_ARGUMENTS (decl) = t;
+ }
+
+ tree data_name = get_identifier (".omp_data_i");
+ t = build_decl (DECL_SOURCE_LOCATION (decl), PARM_DECL, data_name,
+ ptr_type_node);
DECL_ARTIFICIAL (t) = 1;
DECL_NAMELESS (t) = 1;
DECL_ARG_TYPE (t) = ptr_type_node;
DECL_CONTEXT (t) = current_function_decl;
TREE_USED (t) = 1;
+ if (cilk_for_count)
+ DECL_CHAIN (t) = DECL_ARGUMENTS (decl);
DECL_ARGUMENTS (decl) = t;
if (!task_copy)
ctx->receiver_decl = t;
{
omp_context *ctx;
tree name;
- gimple stmt = gsi_stmt (*gsi);
+ gomp_parallel *stmt = as_a <gomp_parallel *> (gsi_stmt (*gsi));
/* Ignore parallel directives with empty bodies, unless there
are copyin clauses. */
if (gimple_omp_parallel_combined_p (stmt))
{
- gimple for_stmt;
struct walk_stmt_info wi;
memset (&wi, 0, sizeof (wi));
wi.val_only = true;
walk_gimple_seq (gimple_omp_body (stmt),
find_combined_for, NULL, &wi);
- for_stmt = (gimple) wi.info;
- if (for_stmt)
+ if (wi.info)
{
+ gomp_for *for_stmt = as_a <gomp_for *> ((gimple) wi.info);
struct omp_for_data fd;
extract_omp_for_data (for_stmt, &fd, NULL);
/* We need two temporaries with fd.loop.v type (istart/iend)
count += fd.collapse - 1;
for (i = 0; i < count; i++)
{
- tree temp = create_tmp_var (type, NULL);
+ tree temp = create_tmp_var (type);
tree c = build_omp_clause (UNKNOWN_LOCATION,
OMP_CLAUSE__LOOPTEMP_);
+ insert_decl_map (&outer_ctx->cb, temp, temp);
OMP_CLAUSE_DECL (c) = temp;
OMP_CLAUSE_CHAIN (c) = gimple_omp_parallel_clauses (stmt);
gimple_omp_parallel_set_clauses (stmt, c);
}
ctx = new_omp_context (stmt, outer_ctx);
+ taskreg_contexts.safe_push (ctx);
if (taskreg_nesting_level > 1)
ctx->is_nested = true;
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
DECL_ARTIFICIAL (name) = 1;
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
+ TYPE_ARTIFICIAL (ctx->record_type) = 1;
create_omp_child_function (ctx, false);
gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn);
if (TYPE_FIELDS (ctx->record_type) == NULL)
ctx->record_type = ctx->receiver_decl = NULL;
- else
- {
- layout_type (ctx->record_type);
- fixup_child_record_type (ctx);
- }
}
/* Scan an OpenMP task directive. */
{
omp_context *ctx;
tree name, t;
- gimple stmt = gsi_stmt (*gsi);
- location_t loc = gimple_location (stmt);
+ gomp_task *stmt = as_a <gomp_task *> (gsi_stmt (*gsi));
/* Ignore task directives with empty bodies. */
if (optimize > 0
}
ctx = new_omp_context (stmt, outer_ctx);
+ taskreg_contexts.safe_push (ctx);
if (taskreg_nesting_level > 1)
ctx->is_nested = true;
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
DECL_ARTIFICIAL (name) = 1;
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
+ TYPE_ARTIFICIAL (ctx->record_type) = 1;
create_omp_child_function (ctx, false);
gimple_omp_task_set_child_fn (stmt, ctx->cb.dst_fn);
DECL_ARTIFICIAL (name) = 1;
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->srecord_type) = name;
+ TYPE_ARTIFICIAL (ctx->srecord_type) = 1;
create_omp_child_function (ctx, true);
}
t = build_int_cst (long_integer_type_node, 1);
gimple_omp_task_set_arg_align (stmt, t);
}
+}
+
+
+/* If any decls have been made addressable during scan_omp,
+ adjust their fields if needed, and layout record types
+ of parallel/task constructs. */
+
+static void
+finish_taskreg_scan (omp_context *ctx)
+{
+ if (ctx->record_type == NULL_TREE)
+ return;
+
+ /* If any task_shared_vars were needed, verify all
+ OMP_CLAUSE_SHARED clauses on GIMPLE_OMP_{PARALLEL,TASK}
+ statements if use_pointer_for_field hasn't changed
+ because of that. If it did, update field types now. */
+ if (task_shared_vars)
+ {
+ tree c;
+
+ for (c = gimple_omp_taskreg_clauses (ctx->stmt);
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED)
+ {
+ tree decl = OMP_CLAUSE_DECL (c);
+
+ /* Global variables don't need to be copied,
+ the receiver side will use them directly. */
+ if (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
+ continue;
+ if (!bitmap_bit_p (task_shared_vars, DECL_UID (decl))
+ || !use_pointer_for_field (decl, ctx))
+ continue;
+ tree field = lookup_field (decl, ctx);
+ if (TREE_CODE (TREE_TYPE (field)) == POINTER_TYPE
+ && TREE_TYPE (TREE_TYPE (field)) == TREE_TYPE (decl))
+ continue;
+ TREE_TYPE (field) = build_pointer_type (TREE_TYPE (decl));
+ TREE_THIS_VOLATILE (field) = 0;
+ DECL_USER_ALIGN (field) = 0;
+ DECL_ALIGN (field) = TYPE_ALIGN (TREE_TYPE (field));
+ if (TYPE_ALIGN (ctx->record_type) < DECL_ALIGN (field))
+ TYPE_ALIGN (ctx->record_type) = DECL_ALIGN (field);
+ if (ctx->srecord_type)
+ {
+ tree sfield = lookup_sfield (decl, ctx);
+ TREE_TYPE (sfield) = TREE_TYPE (field);
+ TREE_THIS_VOLATILE (sfield) = 0;
+ DECL_USER_ALIGN (sfield) = 0;
+ DECL_ALIGN (sfield) = DECL_ALIGN (field);
+ if (TYPE_ALIGN (ctx->srecord_type) < DECL_ALIGN (sfield))
+ TYPE_ALIGN (ctx->srecord_type) = DECL_ALIGN (sfield);
+ }
+ }
+ }
+
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_PARALLEL)
+ {
+ layout_type (ctx->record_type);
+ fixup_child_record_type (ctx);
+ }
else
{
+ location_t loc = gimple_location (ctx->stmt);
tree *p, vla_fields = NULL_TREE, *q = &vla_fields;
/* Move VLA fields to the end. */
p = &TYPE_FIELDS (ctx->record_type);
fixup_child_record_type (ctx);
if (ctx->srecord_type)
layout_type (ctx->srecord_type);
- t = fold_convert_loc (loc, long_integer_type_node,
- TYPE_SIZE_UNIT (ctx->record_type));
- gimple_omp_task_set_arg_size (stmt, t);
+ tree t = fold_convert_loc (loc, long_integer_type_node,
+ TYPE_SIZE_UNIT (ctx->record_type));
+ gimple_omp_task_set_arg_size (ctx->stmt, t);
t = build_int_cst (long_integer_type_node,
TYPE_ALIGN_UNIT (ctx->record_type));
- gimple_omp_task_set_arg_align (stmt, t);
+ gimple_omp_task_set_arg_align (ctx->stmt, t);
}
}
-/* Scan an OpenMP loop directive. */
+static omp_context *
+enclosing_target_ctx (omp_context *ctx)
+{
+ while (ctx != NULL
+ && gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
+ ctx = ctx->outer;
+ gcc_assert (ctx != NULL);
+ return ctx;
+}
+
+static bool
+oacc_loop_or_target_p (gimple stmt)
+{
+ enum gimple_code outer_type = gimple_code (stmt);
+ return ((outer_type == GIMPLE_OMP_TARGET
+ && ((gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ || (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_KERNELS)))
+ || (outer_type == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP));
+}
+
+/* Scan a GIMPLE_OMP_FOR. */
static void
-scan_omp_for (gimple stmt, omp_context *outer_ctx)
+scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
{
+ enum gimple_code outer_type = GIMPLE_ERROR_MARK;
omp_context *ctx;
size_t i;
+ tree clauses = gimple_omp_for_clauses (stmt);
+
+ if (outer_ctx)
+ outer_type = gimple_code (outer_ctx->stmt);
ctx = new_omp_context (stmt, outer_ctx);
- scan_sharing_clauses (gimple_omp_for_clauses (stmt), ctx);
+ if (is_gimple_omp_oacc (stmt))
+ {
+ if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
+ ctx->gwv_this = outer_ctx->gwv_this;
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ int val;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)
+ val = MASK_GANG;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WORKER)
+ val = MASK_WORKER;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR)
+ val = MASK_VECTOR;
+ else
+ continue;
+ ctx->gwv_this |= val;
+ if (!outer_ctx)
+ {
+ /* Skip; not nested inside a region. */
+ continue;
+ }
+ if (!oacc_loop_or_target_p (outer_ctx->stmt))
+ {
+ /* Skip; not nested inside an OpenACC region. */
+ continue;
+ }
+ if (outer_type == GIMPLE_OMP_FOR)
+ outer_ctx->gwv_below |= val;
+ if (OMP_CLAUSE_OPERAND (c, 0) != NULL_TREE)
+ {
+ omp_context *enclosing = enclosing_target_ctx (outer_ctx);
+ if (gimple_omp_target_kind (enclosing->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ error_at (gimple_location (stmt),
+ "no arguments allowed to gang, worker and vector clauses inside parallel");
+ }
+ }
+ }
+
+ scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_for_pre_body_ptr (stmt), ctx);
for (i = 0; i < gimple_omp_for_collapse (stmt); i++)
scan_omp_op (gimple_omp_for_incr_ptr (stmt, i), ctx);
}
scan_omp (gimple_omp_body_ptr (stmt), ctx);
+
+ if (is_gimple_omp_oacc (stmt))
+ {
+ if (ctx->gwv_this & ctx->gwv_below)
+ error_at (gimple_location (stmt),
+ "gang, worker and vector may occur only once in a loop nest");
+ else if (ctx->gwv_below != 0
+ && ctx->gwv_this > ctx->gwv_below)
+ error_at (gimple_location (stmt),
+ "gang, worker and vector must occur in this order in a loop nest");
+ if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
+ outer_ctx->gwv_below |= ctx->gwv_below;
+ }
}
/* Scan an OpenMP sections directive. */
static void
-scan_omp_sections (gimple stmt, omp_context *outer_ctx)
+scan_omp_sections (gomp_sections *stmt, omp_context *outer_ctx)
{
omp_context *ctx;
/* Scan an OpenMP single directive. */
static void
-scan_omp_single (gimple stmt, omp_context *outer_ctx)
+scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
{
omp_context *ctx;
tree name;
layout_type (ctx->record_type);
}
-/* Scan an OpenMP target{, data, update} directive. */
+/* Scan a GIMPLE_OMP_TARGET. */
static void
-scan_omp_target (gimple stmt, omp_context *outer_ctx)
+scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
{
omp_context *ctx;
tree name;
- int kind = gimple_omp_target_kind (stmt);
+ bool offloaded = is_gimple_omp_offloaded (stmt);
+ tree clauses = gimple_omp_target_clauses (stmt);
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
DECL_ARTIFICIAL (name) = 1;
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ TYPE_ARTIFICIAL (ctx->record_type) = 1;
+ if (offloaded)
{
+ if (is_gimple_omp_oacc (stmt))
+ ctx->reduction_map = splay_tree_new (splay_tree_compare_pointers,
+ 0, 0);
+
create_omp_child_function (ctx, false);
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
- scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
+ if (is_gimple_omp_oacc (stmt))
+ {
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_GANGS)
+ ctx->gwv_this |= MASK_GANG;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_WORKERS)
+ ctx->gwv_this |= MASK_WORKER;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR_LENGTH)
+ ctx->gwv_this |= MASK_VECTOR;
+ }
+ }
+
+ scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
gcc_assert (DECL_ALIGN (field) == align);
#endif
layout_type (ctx->record_type);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
fixup_child_record_type (ctx);
}
}
/* Scan an OpenMP teams directive. */
static void
-scan_omp_teams (gimple stmt, omp_context *outer_ctx)
+scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
{
omp_context *ctx = new_omp_context (stmt, outer_ctx);
scan_sharing_clauses (gimple_omp_teams_clauses (stmt), ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
}
-/* Check OpenMP nesting restrictions. */
+/* Check nesting restrictions. */
static bool
check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
{
+ /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
+ inside an OpenACC CTX. */
+ if (!(is_gimple_omp (stmt)
+ && is_gimple_omp_oacc (stmt)))
+ {
+ for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ if (is_gimple_omp (ctx_->stmt)
+ && is_gimple_omp_oacc (ctx_->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "non-OpenACC construct inside of OpenACC region");
+ return false;
+ }
+ }
+
if (ctx != NULL)
{
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
{
error_at (gimple_location (stmt),
"OpenMP constructs may not be nested inside simd region");
switch (gimple_code (stmt))
{
case GIMPLE_OMP_FOR:
- if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_KIND_SIMD)
+ if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD)
return true;
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
{
}
break;
case GIMPLE_OMP_CRITICAL:
- for (; ctx != NULL; ctx = ctx->outer)
- if (gimple_code (ctx->stmt) == GIMPLE_OMP_CRITICAL
- && (gimple_omp_critical_name (stmt)
- == gimple_omp_critical_name (ctx->stmt)))
- {
- error_at (gimple_location (stmt),
- "critical region may not be nested inside a critical "
- "region with the same name");
- return false;
- }
+ {
+ tree this_stmt_name
+ = gimple_omp_critical_name (as_a <gomp_critical *> (stmt));
+ for (; ctx != NULL; ctx = ctx->outer)
+ if (gomp_critical *other_crit
+ = dyn_cast <gomp_critical *> (ctx->stmt))
+ if (this_stmt_name == gimple_omp_critical_name (other_crit))
+ {
+ error_at (gimple_location (stmt),
+ "critical region may not be nested inside a critical "
+ "region with the same name");
+ return false;
+ }
+ }
break;
case GIMPLE_OMP_TEAMS:
if (ctx == NULL
return false;
}
break;
+ case GIMPLE_OMP_TARGET:
+ for (; ctx != NULL; ctx = ctx->outer)
+ {
+ if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
+ {
+ if (is_gimple_omp (stmt)
+ && is_gimple_omp_oacc (stmt)
+ && is_gimple_omp (ctx->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "OpenACC construct inside of non-OpenACC region");
+ return false;
+ }
+ continue;
+ }
+
+ const char *stmt_name, *ctx_stmt_name;
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION: stmt_name = "target"; break;
+ case GF_OMP_TARGET_KIND_DATA: stmt_name = "target data"; break;
+ case GF_OMP_TARGET_KIND_UPDATE: stmt_name = "target update"; break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break;
+ case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: stmt_name = "enter/exit data"; break;
+ default: gcc_unreachable ();
+ }
+ switch (gimple_omp_target_kind (ctx->stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION: ctx_stmt_name = "target"; break;
+ case GF_OMP_TARGET_KIND_DATA: ctx_stmt_name = "target data"; break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL: ctx_stmt_name = "parallel"; break;
+ case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
+ default: gcc_unreachable ();
+ }
+
+ /* OpenACC/OpenMP mismatch? */
+ if (is_gimple_omp_oacc (stmt)
+ != is_gimple_omp_oacc (ctx->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "%s %s construct inside of %s %s region",
+ (is_gimple_omp_oacc (stmt)
+ ? "OpenACC" : "OpenMP"), stmt_name,
+ (is_gimple_omp_oacc (ctx->stmt)
+ ? "OpenACC" : "OpenMP"), ctx_stmt_name);
+ return false;
+ }
+ if (is_gimple_omp_offloaded (ctx->stmt))
+ {
+ /* No GIMPLE_OMP_TARGET inside offloaded OpenACC CTX. */
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ error_at (gimple_location (stmt),
+ "%s construct inside of %s region",
+ stmt_name, ctx_stmt_name);
+ return false;
+ }
+ else
+ {
+ gcc_checking_assert (!is_gimple_omp_oacc (stmt));
+ warning_at (gimple_location (stmt), 0,
+ "%s construct inside of %s region",
+ stmt_name, ctx_stmt_name);
+ }
+ }
+ }
+ break;
default:
break;
}
/* Helper function scan_omp.
Callback for walk_tree or operators in walk_gimple_stmt used to
- scan for OpenMP directives in TP. */
+ scan for OMP directives in TP. */
static tree
scan_omp_1_op (tree *tp, int *walk_subtrees, void *data)
/* Helper function for scan_omp.
- Callback for walk_gimple_stmt used to scan for OpenMP directives in
+ Callback for walk_gimple_stmt used to scan for OMP directives in
the current statement in GSI. */
static tree
if (gimple_has_location (stmt))
input_location = gimple_location (stmt);
- /* Check the OpenMP nesting restrictions. */
+ /* Check the nesting restrictions. */
bool remove = false;
if (is_gimple_omp (stmt))
remove = !check_omp_nesting_restrictions (stmt, ctx);
if (setjmp_or_longjmp_p (fndecl)
&& ctx
&& gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
{
remove = true;
error_at (gimple_location (stmt),
break;
case GIMPLE_OMP_FOR:
- scan_omp_for (stmt, ctx);
+ scan_omp_for (as_a <gomp_for *> (stmt), ctx);
break;
case GIMPLE_OMP_SECTIONS:
- scan_omp_sections (stmt, ctx);
+ scan_omp_sections (as_a <gomp_sections *> (stmt), ctx);
break;
case GIMPLE_OMP_SINGLE:
- scan_omp_single (stmt, ctx);
+ scan_omp_single (as_a <gomp_single *> (stmt), ctx);
break;
case GIMPLE_OMP_SECTION:
break;
case GIMPLE_OMP_TARGET:
- scan_omp_target (stmt, ctx);
+ scan_omp_target (as_a <gomp_target *> (stmt), ctx);
break;
case GIMPLE_OMP_TEAMS:
- scan_omp_teams (stmt, ctx);
+ scan_omp_teams (as_a <gomp_teams *> (stmt), ctx);
break;
case GIMPLE_BIND:
*handled_ops_p = false;
if (ctx)
- for (var = gimple_bind_vars (stmt); var ; var = DECL_CHAIN (var))
+ for (var = gimple_bind_vars (as_a <gbind *> (stmt));
+ var ;
+ var = DECL_CHAIN (var))
insert_decl_map (&ctx->cb, var, var);
}
break;
/* Scan all the statements starting at the current statement. CTX
- contains context information about the OpenMP directives and
+ contains context information about the OMP directives and
clauses found during the scan. */
static void
{
tree fndecl = builtin_decl_explicit (lhs ? BUILT_IN_GOMP_BARRIER_CANCEL
: BUILT_IN_GOMP_BARRIER);
- gimple g = gimple_build_call (fndecl, 0);
+ gcall *g = gimple_build_call (fndecl, 0);
if (lhs)
gimple_call_set_lhs (g, lhs);
return g;
if (SCALAR_FLOAT_TYPE_P (type))
{
REAL_VALUE_TYPE max, min;
- if (HONOR_INFINITIES (TYPE_MODE (type)))
+ if (HONOR_INFINITIES (type))
{
real_inf (&max);
real_arithmetic (&min, NEGATE_EXPR, &max, NULL);
if (SCALAR_FLOAT_TYPE_P (type))
{
REAL_VALUE_TYPE max;
- if (HONOR_INFINITIES (TYPE_MODE (type)))
+ if (HONOR_INFINITIES (type))
real_inf (&max);
else
real_maxval (&max, 0, TYPE_MODE (type));
/* Otherwise return implementation defined alignment. */
unsigned int al = 1;
- enum machine_mode mode, vmode;
+ machine_mode mode, vmode;
int vs = targetm.vectorize.autovectorize_vector_sizes ();
if (vs)
vs = 1 << floor_log2 (vs);
vs = 1 << floor_log2 (vs);
return vs;
}
- enum machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
+ machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
return GET_MODE_NUNITS (vqimode);
return 1;
{
tree c = find_omp_clause (gimple_omp_for_clauses (ctx->stmt),
OMP_CLAUSE_SAFELEN);
- if (c
- && compare_tree_int (OMP_CLAUSE_SAFELEN_EXPR (c), max_vf) == -1)
+ if (c && TREE_CODE (OMP_CLAUSE_SAFELEN_EXPR (c)) != INTEGER_CST)
+ max_vf = 1;
+ else if (c && compare_tree_int (OMP_CLAUSE_SAFELEN_EXPR (c),
+ max_vf) == -1)
max_vf = tree_to_shwi (OMP_CLAUSE_SAFELEN_EXPR (c));
}
if (max_vf > 1)
{
- idx = create_tmp_var (unsigned_type_node, NULL);
- lane = create_tmp_var (unsigned_type_node, NULL);
+ idx = create_tmp_var (unsigned_type_node);
+ lane = create_tmp_var (unsigned_type_node);
}
}
if (max_vf == 1)
return false;
tree atype = build_array_type_nelts (TREE_TYPE (new_var), max_vf);
- tree avar = create_tmp_var_raw (atype, NULL);
+ tree avar = create_tmp_var_raw (atype);
if (TREE_ADDRESSABLE (new_var))
TREE_ADDRESSABLE (avar) = 1;
DECL_ATTRIBUTES (avar)
return true;
}
+/* Helper function of lower_rec_input_clauses. For a reference
+ in simd reduction, add an underlying variable it will reference. */
+
+static void
+handle_simd_reference (location_t loc, tree new_vard, gimple_seq *ilist)
+{
+ tree z = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_vard)));
+ if (TREE_CONSTANT (z))
+ {
+ const char *name = NULL;
+ if (DECL_NAME (new_vard))
+ name = IDENTIFIER_POINTER (DECL_NAME (new_vard));
+
+ z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)), name);
+ gimple_add_tmp_var (z);
+ TREE_ADDRESSABLE (z) = 1;
+ z = build_fold_addr_expr_loc (loc, z);
+ gimplify_assign (new_vard, z, ilist);
+ }
+}
+
/* Generate code to implement the input clauses, FIRSTPRIVATE and COPYIN,
from the receiver (aka child) side and initializers for REFERENCE_TYPE
private variables. Initialization statements go in ILIST, while calls
bool reduction_omp_orig_ref = false;
int pass;
bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD);
+ && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD);
int max_vf = 0;
tree lane = NULL_TREE, idx = NULL_TREE;
tree ivar = NULL_TREE, lvar = NULL_TREE;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
+ case OMP_CLAUSE_LINEAR:
+ if (OMP_CLAUSE_LINEAR_ARRAY (c))
+ max_vf = 1;
+ /* FALLTHRU */
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_LASTPRIVATE:
- case OMP_CLAUSE_LINEAR:
if (is_variable_sized (OMP_CLAUSE_DECL (c)))
max_vf = 1;
break;
if (pass != 0)
continue;
}
+ /* Even without corresponding firstprivate, if
+ decl is Fortran allocatable, it needs outer var
+ reference. */
+ else if (pass == 0
+ && lang_hooks.decls.omp_private_outer_ref
+ (OMP_CLAUSE_DECL (c)))
+ lastprivate_firstprivate = true;
break;
case OMP_CLAUSE_ALIGNED:
if (pass == 0)
t = build_call_expr_loc (clause_loc, t2, 2, t,
omp_clause_aligned_alignment (c));
t = fold_convert_loc (clause_loc, ptype, t);
- x = create_tmp_var (ptype, NULL);
+ x = create_tmp_var (ptype);
t = build2 (MODIFY_EXPR, ptype, x, t);
gimplify_and_add (t, ilist);
t = build_simple_mem_ref_loc (clause_loc, x);
if (c_kind != OMP_CLAUSE_FIRSTPRIVATE || !is_task_ctx (ctx))
{
- gimple stmt;
+ gcall *stmt;
tree tmp, atmp;
ptr = DECL_VALUE_EXPR (new_var);
/* void *tmp = __builtin_alloca */
atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
stmt = gimple_build_call (atmp, 1, x);
- tmp = create_tmp_var_raw (ptr_type_node, NULL);
+ tmp = create_tmp_var_raw (ptr_type_node);
gimple_add_tmp_var (tmp);
gimple_call_set_lhs (stmt, tmp);
}
else if (TREE_CONSTANT (x))
{
- /* For reduction with placeholder in SIMD loop,
- defer adding the initialization of the reference,
- because if we decide to use SIMD array for it,
- the initilization could cause expansion ICE. */
- if (c_kind == OMP_CLAUSE_REDUCTION
- && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)
- && is_simd)
+ /* For reduction in SIMD loop, defer adding the
+ initialization of the reference, because if we decide
+ to use SIMD array for it, the initilization could cause
+ expansion ICE. */
+ if (c_kind == OMP_CLAUSE_REDUCTION && is_simd)
x = NULL_TREE;
else
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINEAR
&& gimple_omp_for_combined_into_p (ctx->stmt))
{
- tree stept = POINTER_TYPE_P (TREE_TYPE (x))
- ? sizetype : TREE_TYPE (x);
- tree t = fold_convert (stept,
- OMP_CLAUSE_LINEAR_STEP (c));
- tree c = find_omp_clause (clauses,
- OMP_CLAUSE__LOOPTEMP_);
- gcc_assert (c);
- tree l = OMP_CLAUSE_DECL (c);
- if (fd->collapse == 1)
- {
- tree n1 = fd->loop.n1;
- tree step = fd->loop.step;
- tree itype = TREE_TYPE (l);
- if (POINTER_TYPE_P (itype))
- itype = signed_type_for (itype);
- l = fold_build2 (MINUS_EXPR, itype, l, n1);
- if (TYPE_UNSIGNED (itype)
- && fd->loop.cond_code == GT_EXPR)
- l = fold_build2 (TRUNC_DIV_EXPR, itype,
- fold_build1 (NEGATE_EXPR,
- itype, l),
- fold_build1 (NEGATE_EXPR,
- itype, step));
- else
- l = fold_build2 (TRUNC_DIV_EXPR, itype, l, step);
- }
+ tree t = OMP_CLAUSE_LINEAR_STEP (c);
+ tree stept = TREE_TYPE (t);
+ tree ct = find_omp_clause (clauses,
+ OMP_CLAUSE__LOOPTEMP_);
+ gcc_assert (ct);
+ tree l = OMP_CLAUSE_DECL (ct);
+ tree n1 = fd->loop.n1;
+ tree step = fd->loop.step;
+ tree itype = TREE_TYPE (l);
+ if (POINTER_TYPE_P (itype))
+ itype = signed_type_for (itype);
+ l = fold_build2 (MINUS_EXPR, itype, l, n1);
+ if (TYPE_UNSIGNED (itype)
+ && fd->loop.cond_code == GT_EXPR)
+ l = fold_build2 (TRUNC_DIV_EXPR, itype,
+ fold_build1 (NEGATE_EXPR, itype, l),
+ fold_build1 (NEGATE_EXPR,
+ itype, step));
+ else
+ l = fold_build2 (TRUNC_DIV_EXPR, itype, l, step);
t = fold_build2 (MULT_EXPR, stept,
fold_convert (stept, l), t);
+
+ if (OMP_CLAUSE_LINEAR_ARRAY (c))
+ {
+ x = lang_hooks.decls.omp_clause_linear_ctor
+ (c, new_var, x, t);
+ gimplify_and_add (x, ilist);
+ goto do_dtor;
+ }
+
if (POINTER_TYPE_P (TREE_TYPE (x)))
x = fold_build2 (POINTER_PLUS_EXPR,
TREE_TYPE (x), x, t);
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINEAR)
{
- tree iv = create_tmp_var (TREE_TYPE (new_var), NULL);
+ tree iv = create_tmp_var (TREE_TYPE (new_var));
x = lang_hooks.decls.omp_clause_copy_ctor (c, iv, x);
gimplify_and_add (x, ilist);
gimple_stmt_iterator gsi
= gsi_start_1 (gimple_omp_body_ptr (ctx->stmt));
- gimple g
+ gassign *g
= gimple_build_assign (unshare_expr (lvar), iv);
gsi_insert_before_without_update (&gsi, g,
GSI_SAME_STMT);
- tree stept = POINTER_TYPE_P (TREE_TYPE (iv))
- ? sizetype : TREE_TYPE (iv);
- tree t = fold_convert (stept,
- OMP_CLAUSE_LINEAR_STEP (c));
+ tree t = OMP_CLAUSE_LINEAR_STEP (c);
enum tree_code code = PLUS_EXPR;
if (POINTER_TYPE_P (TREE_TYPE (new_var)))
code = POINTER_PLUS_EXPR;
- g = gimple_build_assign_with_ops (code, iv, iv, t);
+ g = gimple_build_assign (iv, code, iv, t);
gsi_insert_before_without_update (&gsi, g,
GSI_SAME_STMT);
break;
But if they aren't used, we need to emit the deferred
initialization now. */
else if (is_reference (var) && is_simd)
- {
- tree z
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_vard)));
- if (TREE_CONSTANT (z))
- {
- const char *name = NULL;
- if (DECL_NAME (var))
- name = IDENTIFIER_POINTER (DECL_NAME (new_vard));
-
- z = create_tmp_var_raw
- (TREE_TYPE (TREE_TYPE (new_vard)), name);
- gimple_add_tmp_var (z);
- TREE_ADDRESSABLE (z) = 1;
- z = build_fold_addr_expr_loc (clause_loc, z);
- gimplify_assign (new_vard, z, ilist);
- }
- }
+ handle_simd_reference (clause_loc, new_vard, ilist);
x = lang_hooks.decls.omp_clause_default_ctor
- (c, new_var, unshare_expr (x));
+ (c, unshare_expr (new_var),
+ build_outer_var_ref (var, ctx));
if (x)
gimplify_and_add (x, ilist);
if (OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c))
if (code == MINUS_EXPR)
code = PLUS_EXPR;
+ tree new_vard = new_var;
+ if (is_simd && is_reference (var))
+ {
+ gcc_assert (TREE_CODE (new_var) == MEM_REF);
+ new_vard = TREE_OPERAND (new_var, 0);
+ gcc_assert (DECL_P (new_vard));
+ }
if (is_simd
&& lower_rec_simd_input_clauses (new_var, ctx, max_vf,
idx, lane, ivar, lvar))
x = build2 (code, TREE_TYPE (ref), ref, ivar);
ref = build_outer_var_ref (var, ctx);
gimplify_assign (ref, x, &llist[1]);
+
+ if (new_vard != new_var)
+ {
+ SET_DECL_VALUE_EXPR (new_vard,
+ build_fold_addr_expr (lvar));
+ DECL_HAS_VALUE_EXPR_P (new_vard) = 1;
+ }
}
else
{
+ if (is_reference (var) && is_simd)
+ handle_simd_reference (clause_loc, new_vard, ilist);
gimplify_assign (new_var, x, ilist);
if (is_simd)
{
OMP_CLAUSE__SIMDUID__DECL (c) = uid;
OMP_CLAUSE_CHAIN (c) = gimple_omp_for_clauses (ctx->stmt);
gimple_omp_for_set_clauses (ctx->stmt, c);
- g = gimple_build_assign_with_ops (INTEGER_CST, lane,
- build_int_cst (unsigned_type_node, 0),
- NULL_TREE);
+ g = gimple_build_assign (lane, INTEGER_CST,
+ build_int_cst (unsigned_type_node, 0));
gimple_seq_add_stmt (ilist, g);
for (int i = 0; i < 2; i++)
if (llist[i])
{
- tree vf = create_tmp_var (unsigned_type_node, NULL);
+ tree vf = create_tmp_var (unsigned_type_node);
g = gimple_build_call_internal (IFN_GOMP_SIMD_VF, 1, uid);
gimple_call_set_lhs (g, vf);
gimple_seq *seq = i == 0 ? ilist : dlist;
gimple_seq_add_stmt (seq, g);
tree t = build_int_cst (unsigned_type_node, 0);
- g = gimple_build_assign_with_ops (INTEGER_CST, idx, t, NULL_TREE);
+ g = gimple_build_assign (idx, INTEGER_CST, t);
gimple_seq_add_stmt (seq, g);
tree body = create_artificial_label (UNKNOWN_LOCATION);
tree header = create_artificial_label (UNKNOWN_LOCATION);
gimple_seq_add_stmt (seq, gimple_build_label (body));
gimple_seq_add_seq (seq, llist[i]);
t = build_int_cst (unsigned_type_node, 1);
- g = gimple_build_assign_with_ops (PLUS_EXPR, idx, idx, t);
+ g = gimple_build_assign (idx, PLUS_EXPR, idx, t);
gimple_seq_add_stmt (seq, g);
gimple_seq_add_stmt (seq, gimple_build_label (header));
g = gimple_build_cond (LT_EXPR, idx, vf, body, end);
tree c = find_omp_clause (gimple_omp_for_clauses (ctx->stmt),
OMP_CLAUSE_SAFELEN);
if (c == NULL_TREE
- || compare_tree_int (OMP_CLAUSE_SAFELEN_EXPR (c),
- max_vf) == 1)
+ || (TREE_CODE (OMP_CLAUSE_SAFELEN_EXPR (c)) == INTEGER_CST
+ && compare_tree_int (OMP_CLAUSE_SAFELEN_EXPR (c),
+ max_vf) == 1))
{
c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_SAFELEN);
OMP_CLAUSE_SAFELEN_EXPR (c) = build_int_cst (integer_type_node,
if (predicate)
{
- gimple stmt;
+ gcond *stmt;
tree label_true, arm1, arm2;
label = create_artificial_label (UNKNOWN_LOCATION);
}
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
{
simduid = find_omp_clause (orig_clauses, OMP_CLAUSE__SIMDUID_);
if (simduid)
{
if (lastlane == NULL)
{
- lastlane = create_tmp_var (unsigned_type_node, NULL);
- gimple g
+ lastlane = create_tmp_var (unsigned_type_node);
+ gcall *g
= gimple_build_call_internal (IFN_GOMP_SIMD_LAST_LANE,
2, simduid,
TREE_OPERAND (val, 1));
gimple_seq_add_stmt (stmt_list, gimple_build_label (label));
}
+static void
+oacc_lower_reduction_var_helper (gimple_seq *stmt_seqp, omp_context *ctx,
+ tree tid, tree var, tree new_var)
+{
+ /* The atomic add at the end of the sum creates unnecessary
+ write contention on accelerators. To work around this,
+ create an array to store the partial reductions. Later, in
+ lower_omp_for (for openacc), the values of array will be
+ combined. */
+
+ tree t = NULL_TREE, array, x;
+ tree type = get_base_type (var);
+ gimple stmt;
+
+ /* Now insert the partial reductions into the array. */
+
+ /* Find the reduction array. */
+
+ tree ptype = build_pointer_type (type);
+
+ t = lookup_oacc_reduction (oacc_get_reduction_array_id (var), ctx);
+ t = build_receiver_ref (t, false, ctx->outer);
+
+ array = create_tmp_var (ptype);
+ gimplify_assign (array, t, stmt_seqp);
+
+ tree ptr = create_tmp_var (TREE_TYPE (array));
+
+ /* Find the reduction array. */
+
+ /* testing a unary conversion. */
+ tree offset = create_tmp_var (sizetype);
+ gimplify_assign (offset, TYPE_SIZE_UNIT (type),
+ stmt_seqp);
+ t = create_tmp_var (sizetype);
+ gimplify_assign (t, unshare_expr (fold_build1 (NOP_EXPR, sizetype, tid)),
+ stmt_seqp);
+ stmt = gimple_build_assign (offset, MULT_EXPR, offset, t);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Offset expression. Does the POINTER_PLUS_EXPR take care
+ of adding sizeof(var) to the array? */
+ ptr = create_tmp_var (ptype);
+ stmt = gimple_build_assign (unshare_expr (ptr), POINTER_PLUS_EXPR, array,
+ offset);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Move the local sum to gfc$sum[i]. */
+ x = unshare_expr (build_simple_mem_ref (ptr));
+ stmt = gimplify_assign (x, new_var, stmt_seqp);
+}
/* Generate code to implement the REDUCTION clauses. */
{
gimple_seq sub_seq = NULL;
gimple stmt;
- tree x, c;
+ tree x, c, tid = NULL_TREE;
int count = 0;
/* SIMD reductions are handled in lower_rec_input_clauses. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
return;
/* First see if there is exactly one reduction clause. Use OMP_ATOMIC
if (count == 0)
return;
+ /* Initialize thread info for OpenACC. */
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ /* Get the current thread id. */
+ tree call = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
+ tid = create_tmp_var (TREE_TYPE (TREE_TYPE (call)));
+ gimple stmt = gimple_build_call (call, 0);
+ gimple_call_set_lhs (stmt, tid);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+ }
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
{
tree var, ref, new_var;
if (code == MINUS_EXPR)
code = PLUS_EXPR;
- if (count == 1)
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ gcc_checking_assert (!OMP_CLAUSE_REDUCTION_PLACEHOLDER (c));
+
+ oacc_lower_reduction_var_helper (stmt_seqp, ctx, tid, var, new_var);
+ }
+ else if (count == 1)
{
tree addr = build_fold_addr_expr_loc (clause_loc, ref);
gimplify_and_add (x, stmt_seqp);
return;
}
-
- if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
+ else if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
{
tree placeholder = OMP_CLAUSE_REDUCTION_PLACEHOLDER (c);
}
}
+ if (is_gimple_omp_oacc (ctx->stmt))
+ return;
+
stmt = gimple_build_call (builtin_decl_explicit (BUILT_IN_GOMP_ATOMIC_START),
0);
gimple_seq_add_stmt (stmt_seqp, stmt);
/* A convenience function to build an empty GIMPLE_COND with just the
condition. */
-static gimple
+static gcond *
gimple_build_cond_empty (tree cond)
{
enum tree_code pred_code;
static void
expand_parallel_call (struct omp_region *region, basic_block bb,
- gimple entry_stmt, vec<tree, va_gc> *ws_args)
+ gomp_parallel *entry_stmt,
+ vec<tree, va_gc> *ws_args)
{
tree t, t1, t2, val, cond, c, clauses, flags;
gimple_stmt_iterator gsi;
(cond != 0) or (cond ? val : 1u). */
if (cond)
{
- gimple_stmt_iterator gsi;
-
cond = gimple_boolify (cond);
if (integer_zerop (val))
edge e, e_then, e_else;
tree tmp_then, tmp_else, tmp_join, tmp_var;
- tmp_var = create_tmp_var (TREE_TYPE (val), NULL);
+ tmp_var = create_tmp_var (TREE_TYPE (val));
if (gimple_in_ssa_p (cfun))
{
- tmp_then = make_ssa_name (tmp_var, NULL);
- tmp_else = make_ssa_name (tmp_var, NULL);
- tmp_join = make_ssa_name (tmp_var, NULL);
+ tmp_then = make_ssa_name (tmp_var);
+ tmp_else = make_ssa_name (tmp_var);
+ tmp_join = make_ssa_name (tmp_var);
}
else
{
tmp_join = tmp_var;
}
- e = split_block (bb, NULL);
+ e = split_block_after_labels (bb);
cond_bb = e->src;
bb = e->dest;
remove_edge (e);
make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
- if (current_loops)
- {
- add_bb_to_loop (then_bb, cond_bb->loop_father);
- add_bb_to_loop (else_bb, cond_bb->loop_father);
- }
+ add_bb_to_loop (then_bb, cond_bb->loop_father);
+ add_bb_to_loop (else_bb, cond_bb->loop_father);
e_then = make_edge (then_bb, bb, EDGE_FALLTHRU);
e_else = make_edge (else_bb, bb, EDGE_FALLTHRU);
if (gimple_in_ssa_p (cfun))
{
- gimple phi = create_phi_node (tmp_join, bb);
+ gphi *phi = create_phi_node (tmp_join, bb);
add_phi_arg (phi, tmp_then, e_then, UNKNOWN_LOCATION);
add_phi_arg (phi, tmp_else, e_else, UNKNOWN_LOCATION);
}
false, GSI_CONTINUE_LINKING);
}
+/* Insert a function call whose name is FUNC_NAME with the information from
+ ENTRY_STMT into the basic_block BB. */
+
+static void
+expand_cilk_for_call (basic_block bb, gomp_parallel *entry_stmt,
+ vec <tree, va_gc> *ws_args)
+{
+ tree t, t1, t2;
+ gimple_stmt_iterator gsi;
+ vec <tree, va_gc> *args;
+
+ gcc_assert (vec_safe_length (ws_args) == 2);
+ tree func_name = (*ws_args)[0];
+ tree grain = (*ws_args)[1];
+
+ tree clauses = gimple_omp_parallel_clauses (entry_stmt);
+ tree count = find_omp_clause (clauses, OMP_CLAUSE__CILK_FOR_COUNT_);
+ gcc_assert (count != NULL_TREE);
+ count = OMP_CLAUSE_OPERAND (count, 0);
+
+ gsi = gsi_last_bb (bb);
+ t = gimple_omp_parallel_data_arg (entry_stmt);
+ if (t == NULL)
+ t1 = null_pointer_node;
+ else
+ t1 = build_fold_addr_expr (t);
+ t2 = build_fold_addr_expr (gimple_omp_parallel_child_fn (entry_stmt));
+
+ vec_alloc (args, 4);
+ args->quick_push (t2);
+ args->quick_push (t1);
+ args->quick_push (count);
+ args->quick_push (grain);
+ t = build_call_expr_loc_vec (UNKNOWN_LOCATION, func_name, args);
+
+ force_gimple_operand_gsi (&gsi, t, true, NULL_TREE, false,
+ GSI_CONTINUE_LINKING);
+}
/* Build the function call to GOMP_task to actually
generate the task operation. BB is the block where to insert the code. */
static void
-expand_task_call (basic_block bb, gimple entry_stmt)
+expand_task_call (basic_block bb, gomp_task *entry_stmt)
{
tree t, t1, t2, t3, flags, cond, c, c2, clauses, depend;
gimple_stmt_iterator gsi;
of such a variable. */
if (any_addressable_vars < 0)
{
- gimple parallel_stmt = last_stmt (region->entry);
+ gomp_parallel *parallel_stmt
+ = as_a <gomp_parallel *> (last_stmt (region->entry));
tree child_fun = gimple_omp_parallel_child_fn (parallel_stmt);
tree local_decls, block, decl;
unsigned ix;
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
entry_bb = region->entry;
- exit_bb = region->exit;
+ if (gimple_code (entry_stmt) == GIMPLE_OMP_TASK)
+ exit_bb = region->cont;
+ else
+ exit_bb = region->exit;
- if (is_combined_parallel (region))
+ bool is_cilk_for
+ = (flag_cilkplus
+ && gimple_code (entry_stmt) == GIMPLE_OMP_PARALLEL
+ && find_omp_clause (gimple_omp_parallel_clauses (entry_stmt),
+ OMP_CLAUSE__CILK_FOR_COUNT_) != NULL_TREE);
+
+ if (is_cilk_for)
+ /* If it is a _Cilk_for statement, it is modelled *like* a parallel for,
+ and the inner statement contains the name of the built-in function
+ and grain. */
+ ws_args = region->inner->ws_args;
+ else if (is_combined_parallel (region))
ws_args = region->ws_args;
else
ws_args = NULL;
the region, in which case all we need to do is make the
sub-graph unreachable and emit the parallel call. */
edge entry_succ_e, exit_succ_e;
- gimple_stmt_iterator gsi;
entry_succ_e = single_succ_edge (entry_bb);
variable. In which case, we need to keep the assignment. */
if (gimple_omp_taskreg_data_arg (entry_stmt))
{
- basic_block entry_succ_bb = single_succ (entry_bb);
- gimple_stmt_iterator gsi;
+ basic_block entry_succ_bb
+ = single_succ_p (entry_bb) ? single_succ (entry_bb)
+ : FALLTHRU_EDGE (entry_bb)->dest;
tree arg, narg;
gimple parcopy_stmt = NULL;
if (TREE_CODE (t) == VAR_DECL
&& TREE_STATIC (t)
&& !DECL_EXTERNAL (t))
- varpool_finalize_decl (t);
+ varpool_node::finalize_decl (t);
DECL_SAVED_TREE (child_fn) = NULL;
/* We'll create a CFG for child_fn, so no gimple body is needed. */
gimple_set_body (child_fn, NULL);
stmt = gsi_stmt (gsi);
gcc_assert (stmt && (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
|| gimple_code (stmt) == GIMPLE_OMP_TASK));
- gsi_remove (&gsi, true);
e = split_block (entry_bb, stmt);
+ gsi_remove (&gsi, true);
entry_bb = e->dest;
- single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+ edge e2 = NULL;
+ if (gimple_code (entry_stmt) == GIMPLE_OMP_PARALLEL)
+ single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+ else
+ {
+ e2 = make_edge (e->src, BRANCH_EDGE (entry_bb)->dest, EDGE_ABNORMAL);
+ gcc_assert (e2->dest == region->exit);
+ remove_edge (BRANCH_EDGE (entry_bb));
+ set_immediate_dominator (CDI_DOMINATORS, e2->dest, e->src);
+ gsi = gsi_last_bb (region->exit);
+ gcc_assert (!gsi_end_p (gsi)
+ && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+ gsi_remove (&gsi, true);
+ }
- /* Convert GIMPLE_OMP_RETURN into a RETURN_EXPR. */
+ /* Convert GIMPLE_OMP_{RETURN,CONTINUE} into a RETURN_EXPR. */
if (exit_bb)
{
gsi = gsi_last_bb (exit_bb);
gcc_assert (!gsi_end_p (gsi)
- && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+ && (gimple_code (gsi_stmt (gsi))
+ == (e2 ? GIMPLE_OMP_CONTINUE : GIMPLE_OMP_RETURN)));
stmt = gimple_build_return (NULL);
gsi_insert_after (&gsi, stmt, GSI_SAME_STMT);
gsi_remove (&gsi, true);
new_bb = move_sese_region_to_fn (child_cfun, entry_bb, exit_bb, block);
if (exit_bb)
single_succ_edge (new_bb)->flags = EDGE_FALLTHRU;
+ if (e2)
+ {
+ basic_block dest_bb = e2->dest;
+ if (!exit_bb)
+ make_edge (new_bb, dest_bb, EDGE_FALLTHRU);
+ remove_edge (e2);
+ set_immediate_dominator (CDI_DOMINATORS, dest_bb, new_bb);
+ }
/* When the OMP expansion process cannot guarantee an up-to-date
loop tree arrange for the child function to fixup loops. */
if (loops_state_satisfies_p (LOOPS_NEED_FIXUP))
vec_safe_truncate (child_cfun->local_decls, dstidx);
/* Inform the callgraph about the new function. */
- DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties;
- cgraph_add_new_function (child_fn, true);
+ child_cfun->curr_properties = cfun->curr_properties;
+ child_cfun->has_simduid_loops |= cfun->has_simduid_loops;
+ child_cfun->has_force_vectorize_loops |= cfun->has_force_vectorize_loops;
+ cgraph_node *node = cgraph_node::get_create (child_fn);
+ node->parallelized_function = 1;
+ cgraph_node::add_new_function (child_fn, true);
/* Fix the callgraph edges for child_cfun. Those for cfun will be
fixed in a following pass. */
push_cfun (child_cfun);
if (optimize)
optimize_omp_library_calls (entry_stmt);
- rebuild_cgraph_edges ();
+ cgraph_edge::rebuild_edges ();
/* Some EH regions might become dead, see PR34608. If
pass_cleanup_cfg isn't the first pass to happen with the
}
/* Emit a library call to launch the children threads. */
- if (gimple_code (entry_stmt) == GIMPLE_OMP_PARALLEL)
- expand_parallel_call (region, new_bb, entry_stmt, ws_args);
+ if (is_cilk_for)
+ expand_cilk_for_call (new_bb,
+ as_a <gomp_parallel *> (entry_stmt), ws_args);
+ else if (gimple_code (entry_stmt) == GIMPLE_OMP_PARALLEL)
+ expand_parallel_call (region, new_bb,
+ as_a <gomp_parallel *> (entry_stmt), ws_args);
else
- expand_task_call (new_bb, entry_stmt);
+ expand_task_call (new_bb, as_a <gomp_task *> (entry_stmt));
if (gimple_in_ssa_p (cfun))
update_ssa (TODO_update_ssa_only_virtuals);
}
basic_block &l2_dom_bb)
{
tree t, type = TREE_TYPE (fd->loop.v);
- gimple stmt;
edge e, ne;
int i;
fold_convert (itype, fd->loops[i].n2)))
== NULL_TREE || !integer_onep (t)))
{
+ gcond *cond_stmt;
tree n1, n2;
n1 = fold_convert (itype, unshare_expr (fd->loops[i].n1));
n1 = force_gimple_operand_gsi (gsi, n1, true, NULL_TREE,
n2 = fold_convert (itype, unshare_expr (fd->loops[i].n2));
n2 = force_gimple_operand_gsi (gsi, n2, true, NULL_TREE,
true, GSI_SAME_STMT);
- stmt = gimple_build_cond (fd->loops[i].cond_code, n1, n2,
- NULL_TREE, NULL_TREE);
- gsi_insert_before (gsi, stmt, GSI_SAME_STMT);
- if (walk_tree (gimple_cond_lhs_ptr (stmt),
+ cond_stmt = gimple_build_cond (fd->loops[i].cond_code, n1, n2,
+ NULL_TREE, NULL_TREE);
+ gsi_insert_before (gsi, cond_stmt, GSI_SAME_STMT);
+ if (walk_tree (gimple_cond_lhs_ptr (cond_stmt),
expand_omp_regimplify_p, NULL, NULL)
- || walk_tree (gimple_cond_rhs_ptr (stmt),
+ || walk_tree (gimple_cond_rhs_ptr (cond_stmt),
expand_omp_regimplify_p, NULL, NULL))
{
- *gsi = gsi_for_stmt (stmt);
- gimple_regimplify_operands (stmt, gsi);
+ *gsi = gsi_for_stmt (cond_stmt);
+ gimple_regimplify_operands (cond_stmt, gsi);
}
- e = split_block (entry_bb, stmt);
+ e = split_block (entry_bb, cond_stmt);
if (zero_iter_bb == NULL)
{
+ gassign *assign_stmt;
first_zero_iter = i;
zero_iter_bb = create_empty_bb (entry_bb);
- if (current_loops)
- add_bb_to_loop (zero_iter_bb, entry_bb->loop_father);
+ add_bb_to_loop (zero_iter_bb, entry_bb->loop_father);
*gsi = gsi_after_labels (zero_iter_bb);
- stmt = gimple_build_assign (fd->loop.n2,
- build_zero_cst (type));
- gsi_insert_before (gsi, stmt, GSI_SAME_STMT);
+ assign_stmt = gimple_build_assign (fd->loop.n2,
+ build_zero_cst (type));
+ gsi_insert_before (gsi, assign_stmt, GSI_SAME_STMT);
set_immediate_dominator (CDI_DOMINATORS, zero_iter_bb,
entry_bb);
}
tree t = fold_convert (TREE_TYPE (tem), counts[i]);
t = force_gimple_operand_gsi (gsi, t, false, NULL_TREE,
false, GSI_CONTINUE_LINKING);
- gimple stmt = gimple_build_assign (tem, t);
+ gassign *stmt = gimple_build_assign (tem, t);
gsi_insert_after (gsi, stmt, GSI_CONTINUE_LINKING);
}
}
tree type = TREE_TYPE (fd->loop.v);
tree tem = create_tmp_reg (type, ".tem");
- gimple stmt = gimple_build_assign (tem, startvar);
+ gassign *stmt = gimple_build_assign (tem, startvar);
gsi_insert_after (gsi, stmt, GSI_CONTINUE_LINKING);
for (i = fd->collapse - 1; i >= 0; i--)
tree vtype = TREE_TYPE (fd->loops[i].v);
bb = create_empty_bb (last_bb);
- if (current_loops)
- add_bb_to_loop (bb, last_bb->loop_father);
+ add_bb_to_loop (bb, last_bb->loop_father);
gsi = gsi_start_bb (bb);
if (i < fd->collapse - 1)
basic_block entry_bb, cont_bb, exit_bb, l0_bb, l1_bb, collapse_bb;
basic_block l2_bb = NULL, l3_bb = NULL;
gimple_stmt_iterator gsi;
- gimple stmt;
+ gassign *assign_stmt;
bool in_combined_parallel = is_combined_parallel (region);
bool broken_loop = region->cont == NULL;
edge e, ne;
DECL_P (startvar)
&& TREE_ADDRESSABLE (startvar),
NULL_TREE, false, GSI_CONTINUE_LINKING);
- stmt = gimple_build_assign (startvar, t);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (startvar, t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
t = iend0;
if (bias)
false, GSI_CONTINUE_LINKING);
if (endvar)
{
- stmt = gimple_build_assign (endvar, iend);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (endvar, iend);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
if (useless_type_conversion_p (TREE_TYPE (fd->loop.v), TREE_TYPE (iend)))
- stmt = gimple_build_assign (fd->loop.v, iend);
+ assign_stmt = gimple_build_assign (fd->loop.v, iend);
else
- stmt = gimple_build_assign_with_ops (NOP_EXPR, fd->loop.v, iend,
- NULL_TREE);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (fd->loop.v, NOP_EXPR, iend);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
}
if (fd->collapse > 1)
expand_omp_for_init_vars (fd, &gsi, counts, inner_stmt, startvar);
/* Code to control the increment and predicate for the sequential
loop goes in the CONT_BB. */
gsi = gsi_last_bb (cont_bb);
- stmt = gsi_stmt (gsi);
- gcc_assert (gimple_code (stmt) == GIMPLE_OMP_CONTINUE);
- vmain = gimple_omp_continue_control_use (stmt);
- vback = gimple_omp_continue_control_def (stmt);
+ gomp_continue *cont_stmt = as_a <gomp_continue *> (gsi_stmt (gsi));
+ gcc_assert (gimple_code (cont_stmt) == GIMPLE_OMP_CONTINUE);
+ vmain = gimple_omp_continue_control_use (cont_stmt);
+ vback = gimple_omp_continue_control_def (cont_stmt);
if (!gimple_omp_for_combined_p (fd->for_stmt))
{
DECL_P (vback)
&& TREE_ADDRESSABLE (vback),
NULL_TREE, true, GSI_SAME_STMT);
- stmt = gimple_build_assign (vback, t);
- gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+ assign_stmt = gimple_build_assign (vback, t);
+ gsi_insert_before (&gsi, assign_stmt, GSI_SAME_STMT);
t = build2 (fd->loop.cond_code, boolean_type_node,
DECL_P (vback) && TREE_ADDRESSABLE (vback) ? t : vback,
iend);
- stmt = gimple_build_cond_empty (t);
- gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+ gcond *cond_stmt = gimple_build_cond_empty (t);
+ gsi_insert_before (&gsi, cond_stmt, GSI_SAME_STMT);
}
/* Remove GIMPLE_OMP_CONTINUE. */
if (TREE_TYPE (t) != boolean_type_node)
t = fold_build2 (NE_EXPR, boolean_type_node,
t, build_int_cst (TREE_TYPE (t), 0));
- stmt = gimple_build_cond_empty (t);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ gcond *cond_stmt = gimple_build_cond_empty (t);
+ gsi_insert_after (&gsi, cond_stmt, GSI_CONTINUE_LINKING);
}
/* Add the loop cleanup function. */
t = builtin_decl_explicit (BUILT_IN_GOMP_LOOP_END_CANCEL);
else
t = builtin_decl_explicit (BUILT_IN_GOMP_LOOP_END);
- stmt = gimple_build_call (t, 0);
+ gcall *call_stmt = gimple_build_call (t, 0);
if (gimple_omp_return_lhs (gsi_stmt (gsi)))
- gimple_call_set_lhs (stmt, gimple_omp_return_lhs (gsi_stmt (gsi)));
- gsi_insert_after (&gsi, stmt, GSI_SAME_STMT);
+ gimple_call_set_lhs (call_stmt, gimple_omp_return_lhs (gsi_stmt (gsi)));
+ gsi_insert_after (&gsi, call_stmt, GSI_SAME_STMT);
gsi_remove (&gsi, true);
/* Connect the new blocks. */
remove_edge (e);
make_edge (cont_bb, l2_bb, EDGE_FALSE_VALUE);
- if (current_loops)
- add_bb_to_loop (l2_bb, cont_bb->loop_father);
+ add_bb_to_loop (l2_bb, cont_bb->loop_father);
e = find_edge (cont_bb, l1_bb);
if (gimple_omp_for_combined_p (fd->for_stmt))
{
basic_block body_bb, cont_bb, collapse_bb = NULL;
basic_block fin_bb;
gimple_stmt_iterator gsi;
- gimple stmt;
edge ep;
- enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS;
- enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM;
bool broken_loop = region->cont == NULL;
tree *counts = NULL;
tree n1, n2, step;
+ gcc_checking_assert ((gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP)
+ || !inner_stmt);
+
itype = type = TREE_TYPE (fd->loop.v);
if (POINTER_TYPE_P (type))
itype = signed_type_for (type);
gsi = gsi_last_bb (entry_bb);
gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
- if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
- {
- get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS;
- get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM;
- }
-
if (fd->collapse > 1)
{
int first_zero_iter = -1;
n2 = fold_convert (type, unshare_expr (fd->loop.n2));
n2 = force_gimple_operand_gsi (&gsi, n2, true, NULL_TREE,
true, GSI_SAME_STMT);
- stmt = gimple_build_cond (fd->loop.cond_code, n1, n2,
- NULL_TREE, NULL_TREE);
- gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
- if (walk_tree (gimple_cond_lhs_ptr (stmt),
+ gcond *cond_stmt = gimple_build_cond (fd->loop.cond_code, n1, n2,
+ NULL_TREE, NULL_TREE);
+ gsi_insert_before (&gsi, cond_stmt, GSI_SAME_STMT);
+ if (walk_tree (gimple_cond_lhs_ptr (cond_stmt),
expand_omp_regimplify_p, NULL, NULL)
- || walk_tree (gimple_cond_rhs_ptr (stmt),
+ || walk_tree (gimple_cond_rhs_ptr (cond_stmt),
expand_omp_regimplify_p, NULL, NULL))
{
- gsi = gsi_for_stmt (stmt);
- gimple_regimplify_operands (stmt, &gsi);
+ gsi = gsi_for_stmt (cond_stmt);
+ gimple_regimplify_operands (cond_stmt, &gsi);
}
- ep = split_block (entry_bb, stmt);
+ ep = split_block (entry_bb, cond_stmt);
ep->flags = EDGE_TRUE_VALUE;
entry_bb = ep->dest;
ep->probability = REG_BR_PROB_BASE - (REG_BR_PROB_BASE / 2000 - 1);
if (gimple_in_ssa_p (cfun))
{
int dest_idx = find_edge (entry_bb, fin_bb)->dest_idx;
- for (gsi = gsi_start_phis (fin_bb);
- !gsi_end_p (gsi); gsi_next (&gsi))
+ for (gphi_iterator gpi = gsi_start_phis (fin_bb);
+ !gsi_end_p (gpi); gsi_next (&gpi))
{
- gimple phi = gsi_stmt (gsi);
+ gphi *phi = gpi.phi ();
add_phi_arg (phi, gimple_phi_arg_def (phi, dest_idx),
ep, UNKNOWN_LOCATION);
}
gsi = gsi_last_bb (entry_bb);
}
- t = build_call_expr (builtin_decl_explicit (get_num_threads), 0);
- t = fold_convert (itype, t);
- nthreads = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ switch (gimple_omp_for_kind (fd->for_stmt))
+ {
+ case GF_OMP_FOR_KIND_FOR:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
+ break;
+ case GF_OMP_FOR_KIND_DISTRIBUTE:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
+ break;
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ nthreads = builtin_decl_explicit (BUILT_IN_GOACC_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ nthreads = build_call_expr (nthreads, 0);
+ nthreads = fold_convert (itype, nthreads);
+ nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
true, GSI_SAME_STMT);
-
- t = build_call_expr (builtin_decl_explicit (get_thread_num), 0);
- t = fold_convert (itype, t);
- threadid = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ threadid = build_call_expr (threadid, 0);
+ threadid = fold_convert (itype, threadid);
+ threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
true, GSI_SAME_STMT);
n1 = fd->loop.n1;
gsi_insert_before (&gsi, gimple_build_assign (tt, t), GSI_SAME_STMT);
t = build2 (LT_EXPR, boolean_type_node, threadid, tt);
- stmt = gimple_build_cond_empty (t);
- gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+ gcond *cond_stmt = gimple_build_cond_empty (t);
+ gsi_insert_before (&gsi, cond_stmt, GSI_SAME_STMT);
- second_bb = split_block (entry_bb, stmt)->dest;
+ second_bb = split_block (entry_bb, cond_stmt)->dest;
gsi = gsi_last_bb (second_bb);
gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
gsi_insert_before (&gsi, gimple_build_assign (tt, build_int_cst (itype, 0)),
GSI_SAME_STMT);
- stmt = gimple_build_assign_with_ops (PLUS_EXPR, q, q,
- build_int_cst (itype, 1));
- gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+ gassign *assign_stmt
+ = gimple_build_assign (q, PLUS_EXPR, q, build_int_cst (itype, 1));
+ gsi_insert_before (&gsi, assign_stmt, GSI_SAME_STMT);
- third_bb = split_block (second_bb, stmt)->dest;
+ third_bb = split_block (second_bb, assign_stmt)->dest;
gsi = gsi_last_bb (third_bb);
gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
DECL_P (startvar)
&& TREE_ADDRESSABLE (startvar),
NULL_TREE, false, GSI_CONTINUE_LINKING);
- stmt = gimple_build_assign (startvar, t);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (startvar, t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
t = fold_convert (itype, e0);
t = fold_build2 (MULT_EXPR, itype, t, step);
false, GSI_CONTINUE_LINKING);
if (endvar)
{
- stmt = gimple_build_assign (endvar, e);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (endvar, e);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
if (useless_type_conversion_p (TREE_TYPE (fd->loop.v), TREE_TYPE (e)))
- stmt = gimple_build_assign (fd->loop.v, e);
+ assign_stmt = gimple_build_assign (fd->loop.v, e);
else
- stmt = gimple_build_assign_with_ops (NOP_EXPR, fd->loop.v, e,
- NULL_TREE);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (fd->loop.v, NOP_EXPR, e);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
}
if (fd->collapse > 1)
expand_omp_for_init_vars (fd, &gsi, counts, inner_stmt, startvar);
/* The code controlling the sequential loop replaces the
GIMPLE_OMP_CONTINUE. */
gsi = gsi_last_bb (cont_bb);
- stmt = gsi_stmt (gsi);
- gcc_assert (gimple_code (stmt) == GIMPLE_OMP_CONTINUE);
- vmain = gimple_omp_continue_control_use (stmt);
- vback = gimple_omp_continue_control_def (stmt);
+ gomp_continue *cont_stmt = as_a <gomp_continue *> (gsi_stmt (gsi));
+ gcc_assert (gimple_code (cont_stmt) == GIMPLE_OMP_CONTINUE);
+ vmain = gimple_omp_continue_control_use (cont_stmt);
+ vback = gimple_omp_continue_control_def (cont_stmt);
if (!gimple_omp_for_combined_p (fd->for_stmt))
{
DECL_P (vback)
&& TREE_ADDRESSABLE (vback),
NULL_TREE, true, GSI_SAME_STMT);
- stmt = gimple_build_assign (vback, t);
- gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+ assign_stmt = gimple_build_assign (vback, t);
+ gsi_insert_before (&gsi, assign_stmt, GSI_SAME_STMT);
t = build2 (fd->loop.cond_code, boolean_type_node,
DECL_P (vback) && TREE_ADDRESSABLE (vback)
if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
{
t = gimple_omp_return_lhs (gsi_stmt (gsi));
- gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
+ if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ gcc_checking_assert (t == NULL_TREE);
+ else
+ gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
}
gsi_remove (&gsi, true);
{
tree n, s0, e0, e, t;
tree trip_var, trip_init, trip_main, trip_back, nthreads, threadid;
- tree type, itype, v_main, v_back, v_extra;
+ tree type, itype, vmain, vback, vextra;
basic_block entry_bb, exit_bb, body_bb, seq_start_bb, iter_part_bb;
basic_block trip_update_bb = NULL, cont_bb, collapse_bb = NULL, fin_bb;
- gimple_stmt_iterator si;
- gimple stmt;
+ gimple_stmt_iterator gsi;
edge se;
- enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS;
- enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM;
bool broken_loop = region->cont == NULL;
tree *counts = NULL;
tree n1, n2, step;
+ gcc_checking_assert ((gimple_omp_for_kind (fd->for_stmt)
+ != GF_OMP_FOR_KIND_OACC_LOOP)
+ || !inner_stmt);
+
itype = type = TREE_TYPE (fd->loop.v);
if (POINTER_TYPE_P (type))
itype = signed_type_for (type);
exit_bb = region->exit;
/* Trip and adjustment setup goes in ENTRY_BB. */
- si = gsi_last_bb (entry_bb);
- gcc_assert (gimple_code (gsi_stmt (si)) == GIMPLE_OMP_FOR);
-
- if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
- {
- get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS;
- get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM;
- }
+ gsi = gsi_last_bb (entry_bb);
+ gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
if (fd->collapse > 1)
{
basic_block l2_dom_bb = NULL;
counts = XALLOCAVEC (tree, fd->collapse);
- expand_omp_for_init_counts (fd, &si, entry_bb, counts,
+ expand_omp_for_init_counts (fd, &gsi, entry_bb, counts,
fin_bb, first_zero_iter,
l2_dom_bb);
t = NULL_TREE;
&& (t == NULL_TREE || !integer_onep (t)))
{
n1 = fold_convert (type, unshare_expr (fd->loop.n1));
- n1 = force_gimple_operand_gsi (&si, n1, true, NULL_TREE,
+ n1 = force_gimple_operand_gsi (&gsi, n1, true, NULL_TREE,
true, GSI_SAME_STMT);
n2 = fold_convert (type, unshare_expr (fd->loop.n2));
- n2 = force_gimple_operand_gsi (&si, n2, true, NULL_TREE,
+ n2 = force_gimple_operand_gsi (&gsi, n2, true, NULL_TREE,
true, GSI_SAME_STMT);
- stmt = gimple_build_cond (fd->loop.cond_code, n1, n2,
- NULL_TREE, NULL_TREE);
- gsi_insert_before (&si, stmt, GSI_SAME_STMT);
- if (walk_tree (gimple_cond_lhs_ptr (stmt),
+ gcond *cond_stmt = gimple_build_cond (fd->loop.cond_code, n1, n2,
+ NULL_TREE, NULL_TREE);
+ gsi_insert_before (&gsi, cond_stmt, GSI_SAME_STMT);
+ if (walk_tree (gimple_cond_lhs_ptr (cond_stmt),
expand_omp_regimplify_p, NULL, NULL)
- || walk_tree (gimple_cond_rhs_ptr (stmt),
+ || walk_tree (gimple_cond_rhs_ptr (cond_stmt),
expand_omp_regimplify_p, NULL, NULL))
{
- si = gsi_for_stmt (stmt);
- gimple_regimplify_operands (stmt, &si);
+ gsi = gsi_for_stmt (cond_stmt);
+ gimple_regimplify_operands (cond_stmt, &gsi);
}
- se = split_block (entry_bb, stmt);
+ se = split_block (entry_bb, cond_stmt);
se->flags = EDGE_TRUE_VALUE;
entry_bb = se->dest;
se->probability = REG_BR_PROB_BASE - (REG_BR_PROB_BASE / 2000 - 1);
if (gimple_in_ssa_p (cfun))
{
int dest_idx = find_edge (entry_bb, fin_bb)->dest_idx;
- for (si = gsi_start_phis (fin_bb);
- !gsi_end_p (si); gsi_next (&si))
+ for (gphi_iterator gpi = gsi_start_phis (fin_bb);
+ !gsi_end_p (gpi); gsi_next (&gpi))
{
- gimple phi = gsi_stmt (si);
+ gphi *phi = gpi.phi ();
add_phi_arg (phi, gimple_phi_arg_def (phi, dest_idx),
se, UNKNOWN_LOCATION);
}
}
- si = gsi_last_bb (entry_bb);
+ gsi = gsi_last_bb (entry_bb);
}
- t = build_call_expr (builtin_decl_explicit (get_num_threads), 0);
- t = fold_convert (itype, t);
- nthreads = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ switch (gimple_omp_for_kind (fd->for_stmt))
+ {
+ case GF_OMP_FOR_KIND_FOR:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM);
+ break;
+ case GF_OMP_FOR_KIND_DISTRIBUTE:
+ nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS);
+ threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM);
+ break;
+ case GF_OMP_FOR_KIND_OACC_LOOP:
+ nthreads = builtin_decl_explicit (BUILT_IN_GOACC_GET_NUM_THREADS);
+ threadid = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ nthreads = build_call_expr (nthreads, 0);
+ nthreads = fold_convert (itype, nthreads);
+ nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
true, GSI_SAME_STMT);
-
- t = build_call_expr (builtin_decl_explicit (get_thread_num), 0);
- t = fold_convert (itype, t);
- threadid = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ threadid = build_call_expr (threadid, 0);
+ threadid = fold_convert (itype, threadid);
+ threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
true, GSI_SAME_STMT);
n1 = fd->loop.n1;
gcc_assert (innerc);
n2 = OMP_CLAUSE_DECL (innerc);
}
- n1 = force_gimple_operand_gsi (&si, fold_convert (type, n1),
+ n1 = force_gimple_operand_gsi (&gsi, fold_convert (type, n1),
true, NULL_TREE, true, GSI_SAME_STMT);
- n2 = force_gimple_operand_gsi (&si, fold_convert (itype, n2),
+ n2 = force_gimple_operand_gsi (&gsi, fold_convert (itype, n2),
true, NULL_TREE, true, GSI_SAME_STMT);
- step = force_gimple_operand_gsi (&si, fold_convert (itype, step),
+ step = force_gimple_operand_gsi (&gsi, fold_convert (itype, step),
true, NULL_TREE, true, GSI_SAME_STMT);
fd->chunk_size
- = force_gimple_operand_gsi (&si, fold_convert (itype, fd->chunk_size),
+ = force_gimple_operand_gsi (&gsi, fold_convert (itype, fd->chunk_size),
true, NULL_TREE, true, GSI_SAME_STMT);
t = build_int_cst (itype, (fd->loop.cond_code == LT_EXPR ? -1 : 1));
else
t = fold_build2 (TRUNC_DIV_EXPR, itype, t, step);
t = fold_convert (itype, t);
- n = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ n = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
true, GSI_SAME_STMT);
trip_var = create_tmp_reg (itype, ".trip");
if (gimple_in_ssa_p (cfun))
{
- trip_init = make_ssa_name (trip_var, NULL);
- trip_main = make_ssa_name (trip_var, NULL);
- trip_back = make_ssa_name (trip_var, NULL);
+ trip_init = make_ssa_name (trip_var);
+ trip_main = make_ssa_name (trip_var);
+ trip_back = make_ssa_name (trip_var);
}
else
{
trip_back = trip_var;
}
- stmt = gimple_build_assign (trip_init, build_int_cst (itype, 0));
- gsi_insert_before (&si, stmt, GSI_SAME_STMT);
+ gassign *assign_stmt
+ = gimple_build_assign (trip_init, build_int_cst (itype, 0));
+ gsi_insert_before (&gsi, assign_stmt, GSI_SAME_STMT);
t = fold_build2 (MULT_EXPR, itype, threadid, fd->chunk_size);
t = fold_build2 (MULT_EXPR, itype, t, step);
t = fold_build_pointer_plus (n1, t);
else
t = fold_build2 (PLUS_EXPR, type, t, n1);
- v_extra = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
- true, GSI_SAME_STMT);
+ vextra = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
+ true, GSI_SAME_STMT);
/* Remove the GIMPLE_OMP_FOR. */
- gsi_remove (&si, true);
+ gsi_remove (&gsi, true);
/* Iteration space partitioning goes in ITER_PART_BB. */
- si = gsi_last_bb (iter_part_bb);
+ gsi = gsi_last_bb (iter_part_bb);
t = fold_build2 (MULT_EXPR, itype, trip_main, nthreads);
t = fold_build2 (PLUS_EXPR, itype, t, threadid);
t = fold_build2 (MULT_EXPR, itype, t, fd->chunk_size);
- s0 = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ s0 = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
false, GSI_CONTINUE_LINKING);
t = fold_build2 (PLUS_EXPR, itype, s0, fd->chunk_size);
t = fold_build2 (MIN_EXPR, itype, t, n);
- e0 = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ e0 = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
false, GSI_CONTINUE_LINKING);
t = build2 (LT_EXPR, boolean_type_node, s0, n);
- gsi_insert_after (&si, gimple_build_cond_empty (t), GSI_CONTINUE_LINKING);
+ gsi_insert_after (&gsi, gimple_build_cond_empty (t), GSI_CONTINUE_LINKING);
/* Setup code for sequential iteration goes in SEQ_START_BB. */
- si = gsi_start_bb (seq_start_bb);
+ gsi = gsi_start_bb (seq_start_bb);
tree startvar = fd->loop.v;
tree endvar = NULL_TREE;
else
t = fold_build2 (PLUS_EXPR, type, t, n1);
t = fold_convert (TREE_TYPE (startvar), t);
- t = force_gimple_operand_gsi (&si, t,
+ t = force_gimple_operand_gsi (&gsi, t,
DECL_P (startvar)
&& TREE_ADDRESSABLE (startvar),
NULL_TREE, false, GSI_CONTINUE_LINKING);
- stmt = gimple_build_assign (startvar, t);
- gsi_insert_after (&si, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (startvar, t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
t = fold_convert (itype, e0);
t = fold_build2 (MULT_EXPR, itype, t, step);
else
t = fold_build2 (PLUS_EXPR, type, t, n1);
t = fold_convert (TREE_TYPE (startvar), t);
- e = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ e = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
false, GSI_CONTINUE_LINKING);
if (endvar)
{
- stmt = gimple_build_assign (endvar, e);
- gsi_insert_after (&si, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (endvar, e);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
if (useless_type_conversion_p (TREE_TYPE (fd->loop.v), TREE_TYPE (e)))
- stmt = gimple_build_assign (fd->loop.v, e);
+ assign_stmt = gimple_build_assign (fd->loop.v, e);
else
- stmt = gimple_build_assign_with_ops (NOP_EXPR, fd->loop.v, e,
- NULL_TREE);
- gsi_insert_after (&si, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (fd->loop.v, NOP_EXPR, e);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
}
if (fd->collapse > 1)
- expand_omp_for_init_vars (fd, &si, counts, inner_stmt, startvar);
+ expand_omp_for_init_vars (fd, &gsi, counts, inner_stmt, startvar);
if (!broken_loop)
{
/* The code controlling the sequential loop goes in CONT_BB,
replacing the GIMPLE_OMP_CONTINUE. */
- si = gsi_last_bb (cont_bb);
- stmt = gsi_stmt (si);
- gcc_assert (gimple_code (stmt) == GIMPLE_OMP_CONTINUE);
- v_main = gimple_omp_continue_control_use (stmt);
- v_back = gimple_omp_continue_control_def (stmt);
+ gsi = gsi_last_bb (cont_bb);
+ gomp_continue *cont_stmt = as_a <gomp_continue *> (gsi_stmt (gsi));
+ vmain = gimple_omp_continue_control_use (cont_stmt);
+ vback = gimple_omp_continue_control_def (cont_stmt);
if (!gimple_omp_for_combined_p (fd->for_stmt))
{
if (POINTER_TYPE_P (type))
- t = fold_build_pointer_plus (v_main, step);
+ t = fold_build_pointer_plus (vmain, step);
else
- t = fold_build2 (PLUS_EXPR, type, v_main, step);
- if (DECL_P (v_back) && TREE_ADDRESSABLE (v_back))
- t = force_gimple_operand_gsi (&si, t, true, NULL_TREE,
+ t = fold_build2 (PLUS_EXPR, type, vmain, step);
+ if (DECL_P (vback) && TREE_ADDRESSABLE (vback))
+ t = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
true, GSI_SAME_STMT);
- stmt = gimple_build_assign (v_back, t);
- gsi_insert_before (&si, stmt, GSI_SAME_STMT);
+ assign_stmt = gimple_build_assign (vback, t);
+ gsi_insert_before (&gsi, assign_stmt, GSI_SAME_STMT);
t = build2 (fd->loop.cond_code, boolean_type_node,
- DECL_P (v_back) && TREE_ADDRESSABLE (v_back)
- ? t : v_back, e);
- gsi_insert_before (&si, gimple_build_cond_empty (t), GSI_SAME_STMT);
+ DECL_P (vback) && TREE_ADDRESSABLE (vback)
+ ? t : vback, e);
+ gsi_insert_before (&gsi, gimple_build_cond_empty (t), GSI_SAME_STMT);
}
/* Remove GIMPLE_OMP_CONTINUE. */
- gsi_remove (&si, true);
+ gsi_remove (&gsi, true);
if (fd->collapse > 1 && !gimple_omp_for_combined_p (fd->for_stmt))
collapse_bb = extract_omp_for_update_vars (fd, cont_bb, body_bb);
/* Trip update code goes into TRIP_UPDATE_BB. */
- si = gsi_start_bb (trip_update_bb);
+ gsi = gsi_start_bb (trip_update_bb);
t = build_int_cst (itype, 1);
t = build2 (PLUS_EXPR, itype, trip_main, t);
- stmt = gimple_build_assign (trip_back, t);
- gsi_insert_after (&si, stmt, GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (trip_back, t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
}
/* Replace the GIMPLE_OMP_RETURN with a barrier, or nothing. */
- si = gsi_last_bb (exit_bb);
- if (!gimple_omp_return_nowait_p (gsi_stmt (si)))
+ gsi = gsi_last_bb (exit_bb);
+ if (!gimple_omp_return_nowait_p (gsi_stmt (gsi)))
{
- t = gimple_omp_return_lhs (gsi_stmt (si));
- gsi_insert_after (&si, build_omp_barrier (t), GSI_SAME_STMT);
+ t = gimple_omp_return_lhs (gsi_stmt (gsi));
+ if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ gcc_checking_assert (t == NULL_TREE);
+ else
+ gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT);
}
- gsi_remove (&si, true);
+ gsi_remove (&gsi, true);
/* Connect the new blocks. */
find_edge (iter_part_bb, seq_start_bb)->flags = EDGE_TRUE_VALUE;
if (gimple_in_ssa_p (cfun))
{
- gimple_stmt_iterator psi;
- gimple phi;
+ gphi_iterator psi;
+ gphi *phi;
edge re, ene;
- edge_var_map_vector *head;
edge_var_map *vm;
size_t i;
appropriate phi nodes in iter_part_bb instead. */
se = single_pred_edge (fin_bb);
re = single_succ_edge (trip_update_bb);
- head = redirect_edge_var_map_vector (re);
+ vec<edge_var_map> *head = redirect_edge_var_map_vector (re);
ene = single_succ_edge (entry_bb);
psi = gsi_start_phis (fin_bb);
for (i = 0; !gsi_end_p (psi) && head->iterate (i, &vm);
gsi_next (&psi), ++i)
{
- gimple nphi;
+ gphi *nphi;
source_location locus;
- phi = gsi_stmt (psi);
+ phi = psi.phi ();
t = gimple_phi_result (phi);
gcc_assert (t == redirect_edge_var_map_result (vm));
nphi = create_phi_node (t, iter_part_bb);
locus = gimple_phi_arg_location_from_edge (phi, se);
/* A special case -- fd->loop.v is not yet computed in
- iter_part_bb, we need to use v_extra instead. */
+ iter_part_bb, we need to use vextra instead. */
if (t == fd->loop.v)
- t = v_extra;
+ t = vextra;
add_phi_arg (nphi, t, ene, locus);
locus = redirect_edge_var_map_location (vm);
add_phi_arg (nphi, redirect_edge_var_map_def (vm), re, locus);
}
- gcc_assert (!gsi_end_p (psi) && i == head->length ());
+ gcc_assert (gsi_end_p (psi) && i == head->length ());
redirect_edge_var_map_clear (re);
while (1)
{
}
}
+/* A subroutine of expand_omp_for. Generate code for _Cilk_for loop.
+ Given parameters:
+ for (V = N1; V cond N2; V += STEP) BODY;
+
+ where COND is "<" or ">" or "!=", we generate pseudocode
+
+ for (ind_var = low; ind_var < high; ind_var++)
+ {
+ V = n1 + (ind_var * STEP)
+
+ <BODY>
+ }
+
+ In the above pseudocode, low and high are function parameters of the
+ child function. In the function below, we are inserting a temp.
+ variable that will be making a call to two OMP functions that will not be
+ found in the body of _Cilk_for (since OMP_FOR cannot be mixed
+ with _Cilk_for). These functions are replaced with low and high
+ by the function that handles taskreg. */
+
+
+static void
+expand_cilk_for (struct omp_region *region, struct omp_for_data *fd)
+{
+ bool broken_loop = region->cont == NULL;
+ basic_block entry_bb = region->entry;
+ basic_block cont_bb = region->cont;
+
+ gcc_assert (EDGE_COUNT (entry_bb->succs) == 2);
+ gcc_assert (broken_loop
+ || BRANCH_EDGE (entry_bb)->dest == FALLTHRU_EDGE (cont_bb)->dest);
+ basic_block l0_bb = FALLTHRU_EDGE (entry_bb)->dest;
+ basic_block l1_bb, l2_bb;
+
+ if (!broken_loop)
+ {
+ gcc_assert (BRANCH_EDGE (cont_bb)->dest == l0_bb);
+ gcc_assert (EDGE_COUNT (cont_bb->succs) == 2);
+ l1_bb = split_block (cont_bb, last_stmt (cont_bb))->dest;
+ l2_bb = BRANCH_EDGE (entry_bb)->dest;
+ }
+ else
+ {
+ BRANCH_EDGE (entry_bb)->flags &= ~EDGE_ABNORMAL;
+ l1_bb = split_edge (BRANCH_EDGE (entry_bb));
+ l2_bb = single_succ (l1_bb);
+ }
+ basic_block exit_bb = region->exit;
+ basic_block l2_dom_bb = NULL;
+
+ gimple_stmt_iterator gsi = gsi_last_bb (entry_bb);
+
+ /* Below statements until the "tree high_val = ..." are pseudo statements
+ used to pass information to be used by expand_omp_taskreg.
+ low_val and high_val will be replaced by the __low and __high
+ parameter from the child function.
+
+ The call_exprs part is a place-holder, it is mainly used
+ to distinctly identify to the top-level part that this is
+ where we should put low and high (reasoning given in header
+ comment). */
+
+ tree child_fndecl
+ = gimple_omp_parallel_child_fn (
+ as_a <gomp_parallel *> (last_stmt (region->outer->entry)));
+ tree t, low_val = NULL_TREE, high_val = NULL_TREE;
+ for (t = DECL_ARGUMENTS (child_fndecl); t; t = TREE_CHAIN (t))
+ {
+ if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (t)), "__high"))
+ high_val = t;
+ else if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (t)), "__low"))
+ low_val = t;
+ }
+ gcc_assert (low_val && high_val);
+
+ tree type = TREE_TYPE (low_val);
+ tree ind_var = create_tmp_reg (type, "__cilk_ind_var");
+ gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR);
+
+ /* Not needed in SSA form right now. */
+ gcc_assert (!gimple_in_ssa_p (cfun));
+ if (l2_dom_bb == NULL)
+ l2_dom_bb = l1_bb;
+
+ tree n1 = low_val;
+ tree n2 = high_val;
+
+ gimple stmt = gimple_build_assign (ind_var, n1);
+
+ /* Replace the GIMPLE_OMP_FOR statement. */
+ gsi_replace (&gsi, stmt, true);
+
+ if (!broken_loop)
+ {
+ /* Code to control the increment goes in the CONT_BB. */
+ gsi = gsi_last_bb (cont_bb);
+ stmt = gsi_stmt (gsi);
+ gcc_assert (gimple_code (stmt) == GIMPLE_OMP_CONTINUE);
+ stmt = gimple_build_assign (ind_var, PLUS_EXPR, ind_var,
+ build_one_cst (type));
+
+ /* Replace GIMPLE_OMP_CONTINUE. */
+ gsi_replace (&gsi, stmt, true);
+ }
+
+ /* Emit the condition in L1_BB. */
+ gsi = gsi_after_labels (l1_bb);
+ t = fold_build2 (MULT_EXPR, TREE_TYPE (fd->loop.step),
+ fold_convert (TREE_TYPE (fd->loop.step), ind_var),
+ fd->loop.step);
+ if (POINTER_TYPE_P (TREE_TYPE (fd->loop.n1)))
+ t = fold_build2 (POINTER_PLUS_EXPR, TREE_TYPE (fd->loop.n1),
+ fd->loop.n1, fold_convert (sizetype, t));
+ else
+ t = fold_build2 (PLUS_EXPR, TREE_TYPE (fd->loop.n1),
+ fd->loop.n1, fold_convert (TREE_TYPE (fd->loop.n1), t));
+ t = fold_convert (TREE_TYPE (fd->loop.v), t);
+ expand_omp_build_assign (&gsi, fd->loop.v, t);
+
+ /* The condition is always '<' since the runtime will fill in the low
+ and high values. */
+ stmt = gimple_build_cond (LT_EXPR, ind_var, n2, NULL_TREE, NULL_TREE);
+ gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+
+ /* Remove GIMPLE_OMP_RETURN. */
+ gsi = gsi_last_bb (exit_bb);
+ gsi_remove (&gsi, true);
+
+ /* Connect the new blocks. */
+ remove_edge (FALLTHRU_EDGE (entry_bb));
+
+ edge e, ne;
+ if (!broken_loop)
+ {
+ remove_edge (BRANCH_EDGE (entry_bb));
+ make_edge (entry_bb, l1_bb, EDGE_FALLTHRU);
+
+ e = BRANCH_EDGE (l1_bb);
+ ne = FALLTHRU_EDGE (l1_bb);
+ e->flags = EDGE_TRUE_VALUE;
+ }
+ else
+ {
+ single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+
+ ne = single_succ_edge (l1_bb);
+ e = make_edge (l1_bb, l0_bb, EDGE_TRUE_VALUE);
+
+ }
+ ne->flags = EDGE_FALSE_VALUE;
+ e->probability = REG_BR_PROB_BASE * 7 / 8;
+ ne->probability = REG_BR_PROB_BASE / 8;
+
+ set_immediate_dominator (CDI_DOMINATORS, l1_bb, entry_bb);
+ set_immediate_dominator (CDI_DOMINATORS, l2_bb, l2_dom_bb);
+ set_immediate_dominator (CDI_DOMINATORS, l0_bb, l1_bb);
+
+ if (!broken_loop)
+ {
+ struct loop *loop = alloc_loop ();
+ loop->header = l1_bb;
+ loop->latch = cont_bb;
+ add_loop (loop, l1_bb->loop_father);
+ loop->safelen = INT_MAX;
+ }
+
+ /* Pick the correct library function based on the precision of the
+ induction variable type. */
+ tree lib_fun = NULL_TREE;
+ if (TYPE_PRECISION (type) == 32)
+ lib_fun = cilk_for_32_fndecl;
+ else if (TYPE_PRECISION (type) == 64)
+ lib_fun = cilk_for_64_fndecl;
+ else
+ gcc_unreachable ();
+
+ gcc_assert (fd->sched_kind == OMP_CLAUSE_SCHEDULE_CILKFOR);
+
+ /* WS_ARGS contains the library function flavor to call:
+ __libcilkrts_cilk_for_64 or __libcilkrts_cilk_for_32), and the
+ user-defined grain value. If the user does not define one, then zero
+ is passed in by the parser. */
+ vec_alloc (region->ws_args, 2);
+ region->ws_args->quick_push (lib_fun);
+ region->ws_args->quick_push (fd->chunk_size);
+}
/* A subroutine of expand_omp_for. Generate code for a simd non-worksharing
loop. Given parameters:
basic_block entry_bb, cont_bb, exit_bb, l0_bb, l1_bb, l2_bb, l2_dom_bb;
gimple_stmt_iterator gsi;
gimple stmt;
+ gcond *cond_stmt;
bool broken_loop = region->cont == NULL;
edge e, ne;
tree *counts = NULL;
t = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
false, GSI_CONTINUE_LINKING);
t = build2 (fd->loop.cond_code, boolean_type_node, fd->loop.v, t);
- stmt = gimple_build_cond_empty (t);
- gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
- if (walk_tree (gimple_cond_lhs_ptr (stmt), expand_omp_regimplify_p,
+ cond_stmt = gimple_build_cond_empty (t);
+ gsi_insert_after (&gsi, cond_stmt, GSI_CONTINUE_LINKING);
+ if (walk_tree (gimple_cond_lhs_ptr (cond_stmt), expand_omp_regimplify_p,
NULL, NULL)
- || walk_tree (gimple_cond_rhs_ptr (stmt), expand_omp_regimplify_p,
+ || walk_tree (gimple_cond_rhs_ptr (cond_stmt), expand_omp_regimplify_p,
NULL, NULL))
{
- gsi = gsi_for_stmt (stmt);
- gimple_regimplify_operands (stmt, &gsi);
+ gsi = gsi_for_stmt (cond_stmt);
+ gimple_regimplify_operands (cond_stmt, &gsi);
}
/* Remove GIMPLE_OMP_RETURN. */
else
{
safelen = OMP_CLAUSE_SAFELEN_EXPR (safelen);
- if (!tree_fits_uhwi_p (safelen)
- || tree_to_uhwi (safelen) > INT_MAX)
+ if (TREE_CODE (safelen) != INTEGER_CST)
+ loop->safelen = 0;
+ else if (!tree_fits_uhwi_p (safelen)
+ || tree_to_uhwi (safelen) > INT_MAX)
loop->safelen = INT_MAX;
else
loop->safelen = tree_to_uhwi (safelen);
cfun->has_force_vectorize_loops = true;
}
}
+ else if (simduid)
+ cfun->has_simduid_loops = true;
}
-/* Expand the OpenMP loop defined by REGION. */
+/* Expand the OMP loop defined by REGION. */
static void
expand_omp_for (struct omp_region *region, gimple inner_stmt)
= (struct omp_for_data_loop *)
alloca (gimple_omp_for_collapse (last_stmt (region->entry))
* sizeof (struct omp_for_data_loop));
- extract_omp_for_data (last_stmt (region->entry), &fd, loops);
+ extract_omp_for_data (as_a <gomp_for *> (last_stmt (region->entry)),
+ &fd, loops);
region->sched_kind = fd.sched_kind;
gcc_assert (EDGE_COUNT (region->entry->succs) == 2);
original loops from being detected. Fix that up. */
loops_state_set (LOOPS_NEED_FIXUP);
- if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_KIND_SIMD)
+ if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_SIMD)
expand_omp_simd (region, &fd);
+ else if (gimple_omp_for_kind (fd.for_stmt) == GF_OMP_FOR_KIND_CILKFOR)
+ expand_cilk_for (region, &fd);
else if (fd.sched_kind == OMP_CLAUSE_SCHEDULE_STATIC
&& !fd.have_ordered)
{
unsigned len;
basic_block entry_bb, l0_bb, l1_bb, l2_bb, default_bb;
gimple_stmt_iterator si, switch_si;
- gimple sections_stmt, stmt, cont;
+ gomp_sections *sections_stmt;
+ gimple stmt;
+ gomp_continue *cont;
edge_iterator ei;
edge e;
struct omp_region *inner;
/* The call to GOMP_sections_start goes in ENTRY_BB, replacing the
GIMPLE_OMP_SECTIONS statement. */
si = gsi_last_bb (entry_bb);
- sections_stmt = gsi_stmt (si);
+ sections_stmt = as_a <gomp_sections *> (gsi_stmt (si));
gcc_assert (gimple_code (sections_stmt) == GIMPLE_OMP_SECTIONS);
vin = gimple_omp_sections_control (sections_stmt);
if (!is_combined_parallel (region))
gcc_assert (gimple_code (gsi_stmt (switch_si)) == GIMPLE_OMP_SECTIONS_SWITCH);
if (exit_reachable)
{
- cont = last_stmt (l1_bb);
+ cont = as_a <gomp_continue *> (last_stmt (l1_bb));
gcc_assert (gimple_code (cont) == GIMPLE_OMP_CONTINUE);
vmain = gimple_omp_continue_control_use (cont);
vnext = gimple_omp_continue_control_def (cont);
t = gimple_block_label (default_bb);
u = build_case_label (NULL, NULL, t);
make_edge (l0_bb, default_bb, 0);
- if (current_loops)
- add_bb_to_loop (default_bb, current_loops->tree_root);
+ add_bb_to_loop (default_bb, current_loops->tree_root);
stmt = gimple_build_switch (vmain, u, label_vec);
gsi_insert_after (&switch_si, stmt, GSI_SAME_STMT);
location_t loc;
gimple stmt;
tree decl, call, type, itype;
- enum machine_mode imode;
+ machine_mode imode;
bool exchange;
gsi = gsi_last_bb (load_bb);
location_t loc;
enum tree_code code;
bool need_old, need_new;
- enum machine_mode imode;
+ machine_mode imode;
bool seq_cst;
/* We expect to find the following sequences:
tree iaddr_val;
iaddr = create_tmp_reg (build_pointer_type_for_mode (itype, ptr_mode,
- true), NULL);
+ true));
iaddr_val
= force_gimple_operand_gsi (&si,
fold_convert (TREE_TYPE (iaddr), addr),
false, NULL_TREE, true, GSI_SAME_STMT);
stmt = gimple_build_assign (iaddr, iaddr_val);
gsi_insert_before (&si, stmt, GSI_SAME_STMT);
- loadedi = create_tmp_var (itype, NULL);
+ loadedi = create_tmp_var (itype);
if (gimple_in_ssa_p (cfun))
- loadedi = make_ssa_name (loadedi, NULL);
+ loadedi = make_ssa_name (loadedi);
}
else
{
gsi2 = gsi_start_bb (loop_header);
if (gimple_in_ssa_p (cfun))
{
- gimple stmt;
+ gassign *stmt;
x = force_gimple_operand_gsi (&gsi2, x, true, NULL_TREE,
true, GSI_SAME_STMT);
stmt = gimple_build_assign (loaded_val, x);
old_vali = loadedi;
else
{
- old_vali = create_tmp_var (TREE_TYPE (loadedi), NULL);
+ old_vali = create_tmp_var (TREE_TYPE (loadedi));
stmt = gimple_build_assign (old_vali, loadedi);
gsi_insert_before (&si, stmt, GSI_SAME_STMT);
tree addr, tree loaded_val, tree stored_val)
{
gimple_stmt_iterator si;
- gimple stmt;
+ gassign *stmt;
tree t;
si = gsi_last_bb (load_bb);
expand_omp_atomic (struct omp_region *region)
{
basic_block load_bb = region->entry, store_bb = region->exit;
- gimple load = last_stmt (load_bb), store = last_stmt (store_bb);
+ gomp_atomic_load *load = as_a <gomp_atomic_load *> (last_stmt (load_bb));
+ gomp_atomic_store *store = as_a <gomp_atomic_store *> (last_stmt (store_bb));
tree loaded_val = gimple_omp_atomic_load_lhs (load);
tree addr = gimple_omp_atomic_load_rhs (load);
tree stored_val = gimple_omp_atomic_store_val (store);
}
-/* Expand the OpenMP target{, data, update} directive starting at REGION. */
+/* Expand the GIMPLE_OMP_TARGET starting at REGION. */
static void
expand_omp_target (struct omp_region *region)
{
basic_block entry_bb, exit_bb, new_bb;
- struct function *child_cfun = NULL;
- tree child_fn = NULL_TREE, block, t;
+ struct function *child_cfun;
+ tree child_fn, block, t;
gimple_stmt_iterator gsi;
- gimple entry_stmt, stmt;
+ gomp_target *entry_stmt;
+ gimple stmt;
edge e;
+ bool offloaded, data_region;
- entry_stmt = last_stmt (region->entry);
+ entry_stmt = as_a <gomp_target *> (last_stmt (region->entry));
new_bb = region->entry;
- int kind = gimple_omp_target_kind (entry_stmt);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+
+ offloaded = is_gimple_omp_offloaded (entry_stmt);
+ switch (gimple_omp_target_kind (entry_stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ data_region = false;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ data_region = true;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ child_fn = NULL_TREE;
+ child_cfun = NULL;
+ if (offloaded)
{
child_fn = gimple_omp_target_child_fn (entry_stmt);
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
}
+ /* Supported by expand_omp_taskreg, but not here. */
+ if (child_cfun != NULL)
+ gcc_checking_assert (!child_cfun->cfg);
+ gcc_checking_assert (!gimple_in_ssa_p (cfun));
+
entry_bb = region->entry;
exit_bb = region->exit;
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
unsigned srcidx, dstidx, num;
- /* If the target region needs data sent from the parent
+ /* If the offloading region needs data sent from the parent
function, then the very first statement (except possible
- tree profile counter updates) of the parallel body
+ tree profile counter updates) of the offloading body
is a copy assignment .OMP_DATA_I = &.OMP_DATA_O. Since
&.OMP_DATA_O is passed as an argument to the child function,
we need to replace it with the argument as seen by the child
function.
In most cases, this will end up being the identity assignment
- .OMP_DATA_I = .OMP_DATA_I. However, if the parallel body had
+ .OMP_DATA_I = .OMP_DATA_I. However, if the offloading body had
a function call that has been inlined, the original PARM_DECL
.OMP_DATA_I may have been converted into a different local
variable. In which case, we need to keep the assignment. */
- if (gimple_omp_target_data_arg (entry_stmt))
+ tree data_arg = gimple_omp_target_data_arg (entry_stmt);
+ if (data_arg)
{
basic_block entry_succ_bb = single_succ (entry_bb);
gimple_stmt_iterator gsi;
tree arg;
gimple tgtcopy_stmt = NULL;
- tree sender
- = TREE_VEC_ELT (gimple_omp_target_data_arg (entry_stmt), 0);
+ tree sender = TREE_VEC_ELT (data_arg, 0);
for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
{
/* Declare local variables needed in CHILD_CFUN. */
block = DECL_INITIAL (child_fn);
BLOCK_VARS (block) = vec2chain (child_cfun->local_decls);
- /* The gimplifier could record temporaries in target block
+ /* The gimplifier could record temporaries in the offloading block
rather than in containing function's local_decls chain,
which would mean cgraph missed finalizing them. Do it now. */
for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t))
if (TREE_CODE (t) == VAR_DECL
&& TREE_STATIC (t)
&& !DECL_EXTERNAL (t))
- varpool_finalize_decl (t);
+ varpool_node::finalize_decl (t);
DECL_SAVED_TREE (child_fn) = NULL;
/* We'll create a CFG for child_fn, so no gimple body is needed. */
gimple_set_body (child_fn, NULL);
for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t))
DECL_CONTEXT (t) = child_fn;
- /* Split ENTRY_BB at GIMPLE_OMP_TARGET,
+ /* Split ENTRY_BB at GIMPLE_*,
so that it can be moved to the child function. */
gsi = gsi_last_bb (entry_bb);
stmt = gsi_stmt (gsi);
- gcc_assert (stmt && gimple_code (stmt) == GIMPLE_OMP_TARGET
- && gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_REGION);
- gsi_remove (&gsi, true);
+ gcc_assert (stmt
+ && gimple_code (stmt) == gimple_code (entry_stmt));
e = split_block (entry_bb, stmt);
+ gsi_remove (&gsi, true);
entry_bb = e->dest;
single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
gsi_remove (&gsi, true);
}
- /* Move the target region into CHILD_CFUN. */
+ /* Move the offloading region into CHILD_CFUN. */
block = gimple_block (entry_stmt);
vec_safe_truncate (child_cfun->local_decls, dstidx);
/* Inform the callgraph about the new function. */
- DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties;
- cgraph_add_new_function (child_fn, true);
+ child_cfun->curr_properties = cfun->curr_properties;
+ child_cfun->has_simduid_loops |= cfun->has_simduid_loops;
+ child_cfun->has_force_vectorize_loops |= cfun->has_force_vectorize_loops;
+ cgraph_node *node = cgraph_node::get_create (child_fn);
+ node->parallelized_function = 1;
+ cgraph_node::add_new_function (child_fn, true);
+
+#ifdef ENABLE_OFFLOADING
+ /* Add the new function to the offload table. */
+ vec_safe_push (offload_funcs, child_fn);
+#endif
/* Fix the callgraph edges for child_cfun. Those for cfun will be
fixed in a following pass. */
push_cfun (child_cfun);
- rebuild_cgraph_edges ();
+ cgraph_edge::rebuild_edges ();
+
+#ifdef ENABLE_OFFLOADING
+ /* Prevent IPA from removing child_fn as unreachable, since there are no
+ refs from the parent function to child_fn in offload LTO mode. */
+ cgraph_node::get (child_fn)->mark_force_output ();
+#endif
/* Some EH regions might become dead, see PR34608. If
pass_cleanup_cfg isn't the first pass to happen with the
pop_cfun ();
}
- /* Emit a library call to launch the target region, or do data
+ /* Emit a library call to launch the offloading region, or do data
transfers. */
tree t1, t2, t3, t4, device, cond, c, clauses;
enum built_in_function start_ix;
location_t clause_loc;
- clauses = gimple_omp_target_clauses (entry_stmt);
+ switch (gimple_omp_target_kind (entry_stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ start_ix = BUILT_IN_GOMP_TARGET;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ start_ix = BUILT_IN_GOMP_TARGET_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ start_ix = BUILT_IN_GOACC_PARALLEL;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ start_ix = BUILT_IN_GOACC_DATA_START;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ start_ix = BUILT_IN_GOACC_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
+ break;
+ default:
+ gcc_unreachable ();
+ }
- if (kind == GF_OMP_TARGET_KIND_REGION)
- start_ix = BUILT_IN_GOMP_TARGET;
- else if (kind == GF_OMP_TARGET_KIND_DATA)
- start_ix = BUILT_IN_GOMP_TARGET_DATA;
- else
- start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ clauses = gimple_omp_target_clauses (entry_stmt);
- /* By default, the value of DEVICE is -1 (let runtime library choose)
- and there is no conditional. */
+ /* By default, the value of DEVICE is GOMP_DEVICE_ICV (let runtime
+ library choose) and there is no conditional. */
cond = NULL_TREE;
- device = build_int_cst (integer_type_node, -1);
+ device = build_int_cst (integer_type_node, GOMP_DEVICE_ICV);
c = find_omp_clause (clauses, OMP_CLAUSE_IF);
if (c)
c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
if (c)
{
+ /* Even if we pass it to all library function calls, it is currently only
+ defined/used for the OpenMP target ones. */
+ gcc_checking_assert (start_ix == BUILT_IN_GOMP_TARGET
+ || start_ix == BUILT_IN_GOMP_TARGET_DATA
+ || start_ix == BUILT_IN_GOMP_TARGET_UPDATE);
+
device = OMP_CLAUSE_DEVICE_ID (c);
clause_loc = OMP_CLAUSE_LOCATION (c);
}
device = fold_convert_loc (clause_loc, integer_type_node, device);
/* If we found the clause 'if (cond)', build
- (cond ? device : -2). */
+ (cond ? device : GOMP_DEVICE_HOST_FALLBACK). */
if (cond)
{
cond = gimple_boolify (cond);
edge e;
tree tmp_var;
- tmp_var = create_tmp_var (TREE_TYPE (device), NULL);
- if (kind != GF_OMP_TARGET_KIND_REGION)
+ tmp_var = create_tmp_var (TREE_TYPE (device));
+ if (offloaded)
+ e = split_block_after_labels (new_bb);
+ else
{
gsi = gsi_last_bb (new_bb);
gsi_prev (&gsi);
e = split_block (new_bb, gsi_stmt (gsi));
}
- else
- e = split_block (new_bb, NULL);
cond_bb = e->src;
new_bb = e->dest;
remove_edge (e);
gsi = gsi_start_bb (else_bb);
stmt = gimple_build_assign (tmp_var,
- build_int_cst (integer_type_node, -2));
+ build_int_cst (integer_type_node,
+ GOMP_DEVICE_HOST_FALLBACK));
gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
- if (current_loops)
- {
- add_bb_to_loop (then_bb, cond_bb->loop_father);
- add_bb_to_loop (else_bb, cond_bb->loop_father);
- }
+ add_bb_to_loop (then_bb, cond_bb->loop_father);
+ add_bb_to_loop (else_bb, cond_bb->loop_father);
make_edge (then_bb, new_bb, EDGE_FALLTHRU);
make_edge (else_bb, new_bb, EDGE_FALLTHRU);
}
gimple g;
- /* FIXME: This will be address of
- extern char __OPENMP_TARGET__[] __attribute__((visibility ("hidden")))
- symbol, as soon as the linker plugin is able to create it for us. */
- tree openmp_target = build_zero_cst (ptr_type_node);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ /* The maximum number used by any start_ix, without varargs. */
+ auto_vec<tree, 11> args;
+ args.quick_push (device);
+ if (offloaded)
+ args.quick_push (build_fold_addr_expr (child_fn));
+ switch (start_ix)
+ {
+ case BUILT_IN_GOMP_TARGET:
+ case BUILT_IN_GOMP_TARGET_DATA:
+ case BUILT_IN_GOMP_TARGET_UPDATE:
+ /* This const void * is part of the current ABI, but we're not actually
+ using it. */
+ args.quick_push (build_zero_cst (ptr_type_node));
+ break;
+ case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOACC_ENTER_EXIT_DATA:
+ case BUILT_IN_GOACC_PARALLEL:
+ case BUILT_IN_GOACC_UPDATE:
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ args.quick_push (t1);
+ args.quick_push (t2);
+ args.quick_push (t3);
+ args.quick_push (t4);
+ switch (start_ix)
{
- tree fnaddr = build_fold_addr_expr (child_fn);
- g = gimple_build_call (builtin_decl_explicit (start_ix), 7,
- device, fnaddr, openmp_target, t1, t2, t3, t4);
+ case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOMP_TARGET:
+ case BUILT_IN_GOMP_TARGET_DATA:
+ case BUILT_IN_GOMP_TARGET_UPDATE:
+ break;
+ case BUILT_IN_GOACC_PARALLEL:
+ {
+ tree t_num_gangs, t_num_workers, t_vector_length;
+
+ /* Default values for num_gangs, num_workers, and vector_length. */
+ t_num_gangs = t_num_workers = t_vector_length
+ = fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node, integer_one_node);
+ /* ..., but if present, use the value specified by the respective
+ clause, making sure that are of the correct type. */
+ c = find_omp_clause (clauses, OMP_CLAUSE_NUM_GANGS);
+ if (c)
+ t_num_gangs = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_NUM_GANGS_EXPR (c));
+ c = find_omp_clause (clauses, OMP_CLAUSE_NUM_WORKERS);
+ if (c)
+ t_num_workers = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_NUM_WORKERS_EXPR (c));
+ c = find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH);
+ if (c)
+ t_vector_length = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_VECTOR_LENGTH_EXPR (c));
+ args.quick_push (t_num_gangs);
+ args.quick_push (t_num_workers);
+ args.quick_push (t_vector_length);
+ }
+ /* FALLTHRU */
+ case BUILT_IN_GOACC_ENTER_EXIT_DATA:
+ case BUILT_IN_GOACC_UPDATE:
+ {
+ tree t_async;
+ int t_wait_idx;
+
+ /* Default values for t_async. */
+ t_async = fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node,
+ build_int_cst (integer_type_node,
+ GOMP_ASYNC_SYNC));
+ /* ..., but if present, use the value specified by the respective
+ clause, making sure that is of the correct type. */
+ c = find_omp_clause (clauses, OMP_CLAUSE_ASYNC);
+ if (c)
+ t_async = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_ASYNC_EXPR (c));
+
+ args.quick_push (t_async);
+ /* Save the index, and... */
+ t_wait_idx = args.length ();
+ /* ... push a default value. */
+ args.quick_push (fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node,
+ integer_zero_node));
+ c = find_omp_clause (clauses, OMP_CLAUSE_WAIT);
+ if (c)
+ {
+ int n = 0;
+
+ for (; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT)
+ {
+ args.safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ integer_type_node,
+ OMP_CLAUSE_WAIT_EXPR (c)));
+ n++;
+ }
+ }
+
+ /* Now that we know the number, replace the default value. */
+ args.ordered_remove (t_wait_idx);
+ args.quick_insert (t_wait_idx,
+ fold_convert_loc (gimple_location (entry_stmt),
+ integer_type_node,
+ build_int_cst (integer_type_node, n)));
+ }
+ }
+ break;
+ default:
+ gcc_unreachable ();
}
- else
- g = gimple_build_call (builtin_decl_explicit (start_ix), 6,
- device, openmp_target, t1, t2, t3, t4);
+
+ g = gimple_build_call_vec (builtin_decl_explicit (start_ix), args);
gimple_set_location (g, gimple_location (entry_stmt));
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
- if (kind != GF_OMP_TARGET_KIND_REGION)
+ if (!offloaded)
{
g = gsi_stmt (gsi);
gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
gsi_remove (&gsi, true);
}
- if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+ if (data_region
+ && region->exit)
{
gsi = gsi_last_bb (region->exit);
g = gsi_stmt (gsi);
region->exit = bb;
parent = parent->outer;
}
-
else if (code == GIMPLE_OMP_CONTINUE)
{
gcc_assert (parent);
{
/* GIMPLE_OMP_SECTIONS_SWITCH is part of
GIMPLE_OMP_SECTIONS, and we do nothing for it. */
- ;
}
- else if (code == GIMPLE_OMP_TARGET
- && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE)
- new_omp_region (bb, code, parent);
else
{
- /* Otherwise, this directive becomes the parent for a new
- region. */
region = new_omp_region (bb, code, parent);
- parent = region;
+ /* Otherwise... */
+ if (code == GIMPLE_OMP_TARGET)
+ {
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ /* ..., other than for those stand-alone directives... */
+ region = NULL;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ }
+ /* ..., this directive becomes the parent for a new region. */
+ if (region)
+ parent = region;
}
}
gcc_assert (root_omp_region != NULL);
}
-/* Expands omp construct (and its subconstructs) starting in HEAD. */
+/* Expands omp construct (and its subconstructs) starting in HEAD. */
+
+void
+omp_expand_local (basic_block head)
+{
+ build_omp_regions_root (head);
+ if (dump_file && (dump_flags & TDF_DETAILS))
+ {
+ fprintf (dump_file, "\nOMP region tree\n\n");
+ dump_omp_region (dump_file, root_omp_region, 0);
+ fprintf (dump_file, "\n");
+ }
+
+ remove_exit_barriers (root_omp_region);
+ expand_omp (root_omp_region);
+
+ free_omp_regions ();
+}
+
+/* Scan the CFG and build a tree of OMP regions. Return the root of
+ the OMP region tree. */
+
+static void
+build_omp_regions (void)
+{
+ gcc_assert (root_omp_region == NULL);
+ calculate_dominance_info (CDI_DOMINATORS);
+ build_omp_regions_1 (ENTRY_BLOCK_PTR_FOR_FN (cfun), NULL, false);
+}
+
+/* Main entry point for expanding OMP-GIMPLE into runtime calls. */
+
+static unsigned int
+execute_expand_omp (void)
+{
+ build_omp_regions ();
+
+ if (!root_omp_region)
+ return 0;
+
+ if (dump_file)
+ {
+ fprintf (dump_file, "\nOMP region tree\n\n");
+ dump_omp_region (dump_file, root_omp_region, 0);
+ fprintf (dump_file, "\n");
+ }
+
+ remove_exit_barriers (root_omp_region);
+
+ expand_omp (root_omp_region);
+
+ cleanup_tree_cfg ();
+
+ free_omp_regions ();
+
+ return 0;
+}
+
+/* OMP expansion -- the default pass, run before creation of SSA form. */
+
+namespace {
+
+const pass_data pass_data_expand_omp =
+{
+ GIMPLE_PASS, /* type */
+ "ompexp", /* name */
+ OPTGROUP_NONE, /* optinfo_flags */
+ TV_NONE, /* tv_id */
+ PROP_gimple_any, /* properties_required */
+ PROP_gimple_eomp, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ 0, /* todo_flags_finish */
+};
+
+class pass_expand_omp : public gimple_opt_pass
+{
+public:
+ pass_expand_omp (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_expand_omp, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual unsigned int execute (function *)
+ {
+ bool gate = ((flag_cilkplus != 0 || flag_openacc != 0 || flag_openmp != 0
+ || flag_openmp_simd != 0)
+ && !seen_error ());
+
+ /* This pass always runs, to provide PROP_gimple_eomp.
+ But often, there is nothing to do. */
+ if (!gate)
+ return 0;
+
+ return execute_expand_omp ();
+ }
+
+}; // class pass_expand_omp
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_expand_omp (gcc::context *ctxt)
+{
+ return new pass_expand_omp (ctxt);
+}
+
+namespace {
+
+const pass_data pass_data_expand_omp_ssa =
+{
+ GIMPLE_PASS, /* type */
+ "ompexpssa", /* name */
+ OPTGROUP_NONE, /* optinfo_flags */
+ TV_NONE, /* tv_id */
+ PROP_cfg | PROP_ssa, /* properties_required */
+ PROP_gimple_eomp, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ TODO_cleanup_cfg | TODO_rebuild_alias, /* todo_flags_finish */
+};
+
+class pass_expand_omp_ssa : public gimple_opt_pass
+{
+public:
+ pass_expand_omp_ssa (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_expand_omp_ssa, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual bool gate (function *fun)
+ {
+ return !(fun->curr_properties & PROP_gimple_eomp);
+ }
+ virtual unsigned int execute (function *) { return execute_expand_omp (); }
+
+}; // class pass_expand_omp_ssa
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_expand_omp_ssa (gcc::context *ctxt)
+{
+ return new pass_expand_omp_ssa (ctxt);
+}
+\f
+/* Routines to lower OMP directives into OMP-GIMPLE. */
+
+/* Helper function to preform, potentially COMPLEX_TYPE, operation and
+ convert it to gimple. */
+static void
+oacc_gimple_assign (tree dest, tree_code op, tree src, gimple_seq *seq)
+{
+ gimple stmt;
+
+ if (TREE_CODE (TREE_TYPE (dest)) != COMPLEX_TYPE)
+ {
+ stmt = gimple_build_assign (dest, op, dest, src);
+ gimple_seq_add_stmt (seq, stmt);
+ return;
+ }
+
+ tree t = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree rdest = fold_build1 (REALPART_EXPR, TREE_TYPE (TREE_TYPE (dest)), dest);
+ gimplify_assign (t, rdest, seq);
+ rdest = t;
+
+ t = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree idest = fold_build1 (IMAGPART_EXPR, TREE_TYPE (TREE_TYPE (dest)), dest);
+ gimplify_assign (t, idest, seq);
+ idest = t;
+
+ t = create_tmp_var (TREE_TYPE (TREE_TYPE (src)));
+ tree rsrc = fold_build1 (REALPART_EXPR, TREE_TYPE (TREE_TYPE (src)), src);
+ gimplify_assign (t, rsrc, seq);
+ rsrc = t;
+
+ t = create_tmp_var (TREE_TYPE (TREE_TYPE (src)));
+ tree isrc = fold_build1 (IMAGPART_EXPR, TREE_TYPE (TREE_TYPE (src)), src);
+ gimplify_assign (t, isrc, seq);
+ isrc = t;
+
+ tree r = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree i = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree result;
+
+ if (op == PLUS_EXPR)
+ {
+ stmt = gimple_build_assign (r, op, rdest, rsrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign (i, op, idest, isrc);
+ gimple_seq_add_stmt (seq, stmt);
+ }
+ else if (op == MULT_EXPR)
+ {
+ /* Let x = a + ib = dest, y = c + id = src.
+ x * y = (ac - bd) + i(ad + bc) */
+ tree ac = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree bd = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree ad = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+ tree bc = create_tmp_var (TREE_TYPE (TREE_TYPE (dest)));
+
+ stmt = gimple_build_assign (ac, MULT_EXPR, rdest, rsrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign (bd, MULT_EXPR, idest, isrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign (r, MINUS_EXPR, ac, bd);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign (ad, MULT_EXPR, rdest, isrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign (bd, MULT_EXPR, idest, rsrc);
+ gimple_seq_add_stmt (seq, stmt);
+
+ stmt = gimple_build_assign (i, PLUS_EXPR, ad, bc);
+ gimple_seq_add_stmt (seq, stmt);
+ }
+ else
+ gcc_unreachable ();
+
+ result = build2 (COMPLEX_EXPR, TREE_TYPE (dest), r, i);
+ gimplify_assign (dest, result, seq);
+}
+
+/* Helper function to initialize local data for the reduction arrays.
+ The reduction arrays need to be placed inside the calling function
+ for accelerators, or else the host won't be able to preform the final
+ reduction. */
-void
-omp_expand_local (basic_block head)
+static void
+oacc_initialize_reduction_data (tree clauses, tree nthreads,
+ gimple_seq *stmt_seqp, omp_context *ctx)
{
- build_omp_regions_root (head);
- if (dump_file && (dump_flags & TDF_DETAILS))
- {
- fprintf (dump_file, "\nOMP region tree\n\n");
- dump_omp_region (dump_file, root_omp_region, 0);
- fprintf (dump_file, "\n");
- }
+ tree c, t, oc;
+ gimple stmt;
+ omp_context *octx;
- remove_exit_barriers (root_omp_region);
- expand_omp (root_omp_region);
+ /* Find the innermost OpenACC parallel context. */
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL))
+ octx = ctx;
+ else
+ octx = ctx->outer;
+ gcc_checking_assert (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (octx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL));
- free_omp_regions ();
+ /* Extract the clauses. */
+ oc = gimple_omp_target_clauses (octx->stmt);
+
+ /* Find the last outer clause. */
+ for (; oc && OMP_CLAUSE_CHAIN (oc); oc = OMP_CLAUSE_CHAIN (oc))
+ ;
+
+ /* Allocate arrays for each reduction variable. */
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
+ continue;
+
+ tree var = OMP_CLAUSE_DECL (c);
+ tree type = get_base_type (var);
+ tree array = lookup_oacc_reduction (oacc_get_reduction_array_id (var),
+ ctx);
+ tree size, call;
+
+ /* Calculate size of the reduction array. */
+ t = create_tmp_var (TREE_TYPE (nthreads));
+ stmt = gimple_build_assign (t, MULT_EXPR, nthreads,
+ fold_convert (TREE_TYPE (nthreads),
+ TYPE_SIZE_UNIT (type)));
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ size = create_tmp_var (sizetype);
+ gimplify_assign (size, fold_build1 (NOP_EXPR, sizetype, t), stmt_seqp);
+
+ /* Now allocate memory for it. */
+ call = unshare_expr (builtin_decl_explicit (BUILT_IN_ALLOCA));
+ stmt = gimple_build_call (call, 1, size);
+ gimple_call_set_lhs (stmt, array);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Map this array into the accelerator. */
+
+ /* Add the reduction array to the list of clauses. */
+ tree x = array;
+ t = build_omp_clause (gimple_location (ctx->stmt), OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (t, GOMP_MAP_FORCE_FROM);
+ OMP_CLAUSE_DECL (t) = x;
+ OMP_CLAUSE_CHAIN (t) = NULL;
+ if (oc)
+ OMP_CLAUSE_CHAIN (oc) = t;
+ else
+ gimple_omp_target_set_clauses (as_a <gomp_target *> (octx->stmt), t);
+ OMP_CLAUSE_SIZE (t) = size;
+ oc = t;
+ }
}
-/* Scan the CFG and build a tree of OMP regions. Return the root of
- the OMP region tree. */
+/* Helper function to process the array of partial reductions. Nthreads
+ indicates the number of threads. Unfortunately, GOACC_GET_NUM_THREADS
+ cannot be used here, because nthreads on the host may be different than
+ on the accelerator. */
static void
-build_omp_regions (void)
+oacc_finalize_reduction_data (tree clauses, tree nthreads,
+ gimple_seq *stmt_seqp, omp_context *ctx)
{
- gcc_assert (root_omp_region == NULL);
- calculate_dominance_info (CDI_DOMINATORS);
- build_omp_regions_1 (ENTRY_BLOCK_PTR_FOR_FN (cfun), NULL, false);
-}
+ tree c, x, var, array, loop_header, loop_body, loop_exit, type;
+ gimple stmt;
-/* Main entry point for expanding OMP-GIMPLE into runtime calls. */
+ /* Create for loop.
-static unsigned int
-execute_expand_omp (void)
-{
- build_omp_regions ();
+ let var = the original reduction variable
+ let array = reduction variable array
- if (!root_omp_region)
- return 0;
+ for (i = 0; i < nthreads; i++)
+ var op= array[i]
+ */
- if (dump_file)
+ loop_header = create_artificial_label (UNKNOWN_LOCATION);
+ loop_body = create_artificial_label (UNKNOWN_LOCATION);
+ loop_exit = create_artificial_label (UNKNOWN_LOCATION);
+
+ /* Create and initialize an index variable. */
+ tree ix = create_tmp_var (sizetype);
+ gimplify_assign (ix, fold_build1 (NOP_EXPR, sizetype, integer_zero_node),
+ stmt_seqp);
+
+ /* Insert the loop header label here. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_label (loop_header));
+
+ /* Exit loop if ix >= nthreads. */
+ x = create_tmp_var (sizetype);
+ gimplify_assign (x, fold_build1 (NOP_EXPR, sizetype, nthreads), stmt_seqp);
+ stmt = gimple_build_cond (GE_EXPR, ix, x, loop_exit, loop_body);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Insert the loop body label here. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_label (loop_body));
+
+ /* Collapse each reduction array, one element at a time. */
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
- fprintf (dump_file, "\nOMP region tree\n\n");
- dump_omp_region (dump_file, root_omp_region, 0);
- fprintf (dump_file, "\n");
- }
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_REDUCTION)
+ continue;
- remove_exit_barriers (root_omp_region);
+ tree_code reduction_code = OMP_CLAUSE_REDUCTION_CODE (c);
- expand_omp (root_omp_region);
+ /* reduction(-:var) sums up the partial results, so it acts
+ identically to reduction(+:var). */
+ if (reduction_code == MINUS_EXPR)
+ reduction_code = PLUS_EXPR;
- cleanup_tree_cfg ();
+ /* Set up reduction variable var. */
+ var = OMP_CLAUSE_DECL (c);
+ type = get_base_type (var);
+ array = lookup_oacc_reduction (oacc_get_reduction_array_id
+ (OMP_CLAUSE_DECL (c)), ctx);
- free_omp_regions ();
+ /* Calculate the array offset. */
+ tree offset = create_tmp_var (sizetype);
+ gimplify_assign (offset, TYPE_SIZE_UNIT (type), stmt_seqp);
+ stmt = gimple_build_assign (offset, MULT_EXPR, offset, ix);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
- return 0;
-}
+ tree ptr = create_tmp_var (TREE_TYPE (array));
+ stmt = gimple_build_assign (ptr, POINTER_PLUS_EXPR, array, offset);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
-/* OMP expansion -- the default pass, run before creation of SSA form. */
+ /* Extract array[ix] into mem. */
+ tree mem = create_tmp_var (type);
+ gimplify_assign (mem, build_simple_mem_ref (ptr), stmt_seqp);
-namespace {
+ /* Find the original reduction variable. */
+ if (is_reference (var))
+ var = build_simple_mem_ref (var);
-const pass_data pass_data_expand_omp =
-{
- GIMPLE_PASS, /* type */
- "ompexp", /* name */
- OPTGROUP_NONE, /* optinfo_flags */
- true, /* has_execute */
- TV_NONE, /* tv_id */
- PROP_gimple_any, /* properties_required */
- 0, /* properties_provided */
- 0, /* properties_destroyed */
- 0, /* todo_flags_start */
- 0, /* todo_flags_finish */
-};
+ tree t = create_tmp_var (type);
-class pass_expand_omp : public gimple_opt_pass
+ x = lang_hooks.decls.omp_clause_assign_op (c, t, var);
+ gimplify_and_add (unshare_expr(x), stmt_seqp);
+
+ /* var = var op mem */
+ switch (OMP_CLAUSE_REDUCTION_CODE (c))
+ {
+ case TRUTH_ANDIF_EXPR:
+ case TRUTH_ORIF_EXPR:
+ t = fold_build2 (OMP_CLAUSE_REDUCTION_CODE (c), integer_type_node,
+ t, mem);
+ gimplify_and_add (t, stmt_seqp);
+ break;
+ default:
+ /* The lhs isn't a gimple_reg when var is COMPLEX_TYPE. */
+ oacc_gimple_assign (t, OMP_CLAUSE_REDUCTION_CODE (c), mem,
+ stmt_seqp);
+ }
+
+ t = fold_build1 (NOP_EXPR, TREE_TYPE (var), t);
+ x = lang_hooks.decls.omp_clause_assign_op (c, var, t);
+ gimplify_and_add (unshare_expr(x), stmt_seqp);
+ }
+
+ /* Increment the induction variable. */
+ tree one = fold_build1 (NOP_EXPR, sizetype, integer_one_node);
+ stmt = gimple_build_assign (ix, PLUS_EXPR, ix, one);
+ gimple_seq_add_stmt (stmt_seqp, stmt);
+
+ /* Go back to the top of the loop. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_goto (loop_header));
+
+ /* Place the loop exit label here. */
+ gimple_seq_add_stmt (stmt_seqp, gimple_build_label (loop_exit));
+}
+
+/* Scan through all of the gimple stmts searching for an OMP_FOR_EXPR, and
+ scan that for reductions. */
+
+static void
+oacc_process_reduction_data (gimple_seq *body, gimple_seq *in_stmt_seqp,
+ gimple_seq *out_stmt_seqp, omp_context *ctx)
{
-public:
- pass_expand_omp (gcc::context *ctxt)
- : gimple_opt_pass (pass_data_expand_omp, ctxt)
- {}
+ gimple_stmt_iterator gsi;
+ gimple_seq inner = NULL;
- /* opt_pass methods: */
- virtual bool gate (function *)
+ /* A collapse clause may have inserted a new bind block. */
+ gsi = gsi_start (*body);
+ while (!gsi_end_p (gsi))
{
- return ((flag_openmp != 0 || flag_openmp_simd != 0
- || flag_cilkplus != 0) && !seen_error ());
+ gimple stmt = gsi_stmt (gsi);
+ if (gbind *bind_stmt = dyn_cast <gbind *> (stmt))
+ {
+ inner = gimple_bind_body (bind_stmt);
+ body = &inner;
+ gsi = gsi_start (*body);
+ }
+ else if (dyn_cast <gomp_for *> (stmt))
+ break;
+ else
+ gsi_next (&gsi);
}
- virtual unsigned int execute (function *) { return execute_expand_omp (); }
+ for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
+ {
+ tree clauses, nthreads, t, c, acc_device, acc_device_host, call,
+ enter, exit;
+ bool reduction_found = false;
-}; // class pass_expand_omp
+ gimple stmt = gsi_stmt (gsi);
-} // anon namespace
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OMP_FOR:
+ clauses = gimple_omp_for_clauses (stmt);
-gimple_opt_pass *
-make_pass_expand_omp (gcc::context *ctxt)
-{
- return new pass_expand_omp (ctxt);
+ /* Search for a reduction clause. */
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ reduction_found = true;
+ break;
+ }
+
+ if (!reduction_found)
+ break;
+
+ ctx = maybe_lookup_ctx (stmt);
+ t = NULL_TREE;
+
+ /* Extract the number of threads. */
+ nthreads = create_tmp_var (sizetype);
+ t = oacc_max_threads (ctx);
+ gimplify_assign (nthreads, t, in_stmt_seqp);
+
+ /* Determine if this is kernel will be executed on the host. */
+ call = builtin_decl_explicit (BUILT_IN_ACC_GET_DEVICE_TYPE);
+ acc_device = create_tmp_var (integer_type_node, ".acc_device_type");
+ stmt = gimple_build_call (call, 0);
+ gimple_call_set_lhs (stmt, acc_device);
+ gimple_seq_add_stmt (in_stmt_seqp, stmt);
+
+ /* Set nthreads = 1 for ACC_DEVICE_TYPE=host. */
+ acc_device_host = create_tmp_var (integer_type_node,
+ ".acc_device_host");
+ gimplify_assign (acc_device_host,
+ build_int_cst (integer_type_node,
+ GOMP_DEVICE_HOST),
+ in_stmt_seqp);
+
+ enter = create_artificial_label (UNKNOWN_LOCATION);
+ exit = create_artificial_label (UNKNOWN_LOCATION);
+
+ stmt = gimple_build_cond (EQ_EXPR, acc_device, acc_device_host,
+ enter, exit);
+ gimple_seq_add_stmt (in_stmt_seqp, stmt);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (enter));
+ gimplify_assign (nthreads, fold_build1 (NOP_EXPR, sizetype,
+ integer_one_node),
+ in_stmt_seqp);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (exit));
+
+ /* Also, set nthreads = 1 for ACC_DEVICE_TYPE=host_nonshm. */
+ gimplify_assign (acc_device_host,
+ build_int_cst (integer_type_node,
+ GOMP_DEVICE_HOST_NONSHM),
+ in_stmt_seqp);
+
+ enter = create_artificial_label (UNKNOWN_LOCATION);
+ exit = create_artificial_label (UNKNOWN_LOCATION);
+
+ stmt = gimple_build_cond (EQ_EXPR, acc_device, acc_device_host,
+ enter, exit);
+ gimple_seq_add_stmt (in_stmt_seqp, stmt);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (enter));
+ gimplify_assign (nthreads, fold_build1 (NOP_EXPR, sizetype,
+ integer_one_node),
+ in_stmt_seqp);
+ gimple_seq_add_stmt (in_stmt_seqp, gimple_build_label (exit));
+
+ oacc_initialize_reduction_data (clauses, nthreads, in_stmt_seqp,
+ ctx);
+ oacc_finalize_reduction_data (clauses, nthreads, out_stmt_seqp, ctx);
+ break;
+ default:
+ // Scan for other directives which support reduction here.
+ break;
+ }
+ }
}
-\f
-/* Routines to lower OpenMP directives into OMP-GIMPLE. */
/* If ctx is a worksharing context inside of a cancellable parallel
region and it isn't nowait, add lhs to its GIMPLE_OMP_RETURN
{
tree fndecl = builtin_decl_explicit (BUILT_IN_GOMP_CANCEL);
tree c_bool_type = TREE_TYPE (TREE_TYPE (fndecl));
- tree lhs = create_tmp_var (c_bool_type, NULL);
+ tree lhs = create_tmp_var (c_bool_type);
gimple_omp_return_set_lhs (omp_return, lhs);
tree fallthru_label = create_artificial_label (UNKNOWN_LOCATION);
gimple g = gimple_build_cond (NE_EXPR, lhs,
{
tree block, control;
gimple_stmt_iterator tgsi;
- gimple stmt, new_stmt, bind, t;
+ gomp_sections *stmt;
+ gimple t;
+ gbind *new_stmt, *bind;
gimple_seq ilist, dlist, olist, new_body;
- stmt = gsi_stmt (*gsi_p);
+ stmt = as_a <gomp_sections *> (gsi_stmt (*gsi_p));
push_gimplify_context ();
to a synchronization analysis pass. */
static void
-lower_omp_single_simple (gimple single_stmt, gimple_seq *pre_p)
+lower_omp_single_simple (gomp_single *single_stmt, gimple_seq *pre_p)
{
location_t loc = gimple_location (single_stmt);
tree tlabel = create_artificial_label (loc);
tree lhs, decl;
decl = builtin_decl_explicit (BUILT_IN_GOMP_SINGLE_START);
- lhs = create_tmp_var (TREE_TYPE (TREE_TYPE (decl)), NULL);
+ lhs = create_tmp_var (TREE_TYPE (TREE_TYPE (decl)));
call = gimple_build_call (decl, 0);
gimple_call_set_lhs (call, lhs);
gimple_seq_add_stmt (pre_p, call);
to a synchronization analysis pass. */
static void
-lower_omp_single_copy (gimple single_stmt, gimple_seq *pre_p, omp_context *ctx)
+lower_omp_single_copy (gomp_single *single_stmt, gimple_seq *pre_p,
+ omp_context *ctx)
{
tree ptr_type, t, l0, l1, l2, bfn_decl;
gimple_seq copyin_seq;
lower_omp_single (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree block;
- gimple t, bind, single_stmt = gsi_stmt (*gsi_p);
+ gimple t;
+ gomp_single *single_stmt = as_a <gomp_single *> (gsi_stmt (*gsi_p));
+ gbind *bind;
gimple_seq bind_body, bind_body_tail = NULL, dlist;
push_gimplify_context ();
lower_omp_master (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree block, lab = NULL, x, bfn_decl;
- gimple stmt = gsi_stmt (*gsi_p), bind;
+ gimple stmt = gsi_stmt (*gsi_p);
+ gbind *bind;
location_t loc = gimple_location (stmt);
gimple_seq tseq;
static void
lower_omp_taskgroup (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
- gimple stmt = gsi_stmt (*gsi_p), bind, x;
+ gimple stmt = gsi_stmt (*gsi_p);
+ gcall *x;
+ gbind *bind;
tree block = make_node (BLOCK);
bind = gimple_build_bind (NULL, NULL, block);
lower_omp_ordered (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree block;
- gimple stmt = gsi_stmt (*gsi_p), bind, x;
+ gimple stmt = gsi_stmt (*gsi_p);
+ gcall *x;
+ gbind *bind;
push_gimplify_context ();
requires that languages coordinate a symbol name. It is therefore
best put here in common code. */
-static GTY((param1_is (tree), param2_is (tree)))
- splay_tree critical_name_mutexes;
+static GTY(()) hash_map<tree, tree> *critical_name_mutexes;
static void
lower_omp_critical (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree block;
tree name, lock, unlock;
- gimple stmt = gsi_stmt (*gsi_p), bind;
+ gomp_critical *stmt = as_a <gomp_critical *> (gsi_stmt (*gsi_p));
+ gbind *bind;
location_t loc = gimple_location (stmt);
gimple_seq tbody;
if (name)
{
tree decl;
- splay_tree_node n;
if (!critical_name_mutexes)
- critical_name_mutexes
- = splay_tree_new_ggc (splay_tree_compare_pointers,
- ggc_alloc_splay_tree_tree_node_tree_node_splay_tree_s,
- ggc_alloc_splay_tree_tree_node_tree_node_splay_tree_node_s);
+ critical_name_mutexes = hash_map<tree, tree>::create_ggc (10);
- n = splay_tree_lookup (critical_name_mutexes, (splay_tree_key) name);
+ tree *n = critical_name_mutexes->get (name);
if (n == NULL)
{
char *new_str;
- decl = create_tmp_var_raw (ptr_type_node, NULL);
+ decl = create_tmp_var_raw (ptr_type_node);
new_str = ACONCAT ((".gomp_critical_user_",
IDENTIFIER_POINTER (name), NULL));
DECL_COMMON (decl) = 1;
DECL_ARTIFICIAL (decl) = 1;
DECL_IGNORED_P (decl) = 1;
- varpool_finalize_decl (decl);
- splay_tree_insert (critical_name_mutexes, (splay_tree_key) name,
- (splay_tree_value) decl);
+ varpool_node::finalize_decl (decl);
+
+ critical_name_mutexes->put (name, decl);
}
else
- decl = (tree) n->value;
+ decl = *n;
+
+ /* If '#pragma omp critical' is inside offloaded region or
+ inside function marked as offloadable, the symbol must be
+ marked as offloadable too. */
+ omp_context *octx;
+ if (cgraph_node::get (current_function_decl)->offloadable)
+ varpool_node::get_create (decl)->offloadable = 1;
+ else
+ for (octx = ctx->outer; octx; octx = octx->outer)
+ if (is_gimple_omp_offloaded (octx->stmt))
+ {
+ varpool_node::get_create (decl)->offloadable = 1;
+ break;
+ }
lock = builtin_decl_explicit (BUILT_IN_GOMP_CRITICAL_NAME_START);
lock = build_call_expr_loc (loc, lock, 1, build_fold_addr_expr_loc (loc, decl));
cond_code = EQ_EXPR;
}
- cond = build2 (cond_code, boolean_type_node, fd->loop.v, fd->loop.n2);
+ tree n2 = fd->loop.n2;
+ if (fd->collapse > 1
+ && TREE_CODE (n2) != INTEGER_CST
+ && gimple_omp_for_combined_into_p (fd->for_stmt)
+ && gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
+ {
+ gomp_for *gfor = as_a <gomp_for *> (ctx->outer->stmt);
+ if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR)
+ {
+ struct omp_for_data outer_fd;
+ extract_omp_for_data (gfor, &outer_fd, NULL);
+ n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+ }
+ }
+ cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2);
clauses = gimple_omp_for_clauses (fd->for_stmt);
stmts = NULL;
}
-/* Lower code for an OpenMP loop directive. */
+/* Lower code for an OMP loop directive. */
static void
lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree *rhs_p, block;
struct omp_for_data fd, *fdp = NULL;
- gimple stmt = gsi_stmt (*gsi_p), new_stmt;
+ gomp_for *stmt = as_a <gomp_for *> (gsi_stmt (*gsi_p));
+ gbind *new_stmt;
gimple_seq omp_for_body, body, dlist;
size_t i;
if (!gimple_seq_empty_p (omp_for_body)
&& gimple_code (gimple_seq_first_stmt (omp_for_body)) == GIMPLE_BIND)
{
- gimple inner_bind = gimple_seq_first_stmt (omp_for_body);
+ gbind *inner_bind
+ = as_a <gbind *> (gimple_seq_first_stmt (omp_for_body));
tree vars = gimple_bind_vars (inner_bind);
gimple_bind_append_vars (new_stmt, vars);
/* bind_vars/BLOCK_VARS are being moved to new_stmt/block, don't
OMP_CLAUSE__LOOPTEMP_);
}
else
- temp = create_tmp_var (type, NULL);
+ {
+ temp = create_tmp_var (type);
+ insert_decl_map (&ctx->outer->cb, temp, temp);
+ }
*pc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__LOOPTEMP_);
OMP_CLAUSE_DECL (*pc) = temp;
pc = &OMP_CLAUSE_CHAIN (*pc);
struct omp_taskcopy_context *tcctx = (struct omp_taskcopy_context *) cb;
if (splay_tree_lookup (tcctx->ctx->sfield_map, (splay_tree_key) var))
- return create_tmp_var (TREE_TYPE (var), NULL);
+ return create_tmp_var (TREE_TYPE (var));
return var;
}
walk_tree (&DECL_FIELD_OFFSET (new_f), copy_tree_body_r,
&tcctx->cb, NULL);
new_fields = new_f;
- *pointer_map_insert (tcctx->cb.decl_map, f) = new_f;
+ tcctx->cb.decl_map->put (f, new_f);
}
TYPE_FIELDS (type) = nreverse (new_fields);
layout_type (type);
/* Create task copyfn. */
static void
-create_task_copyfn (gimple task_stmt, omp_context *ctx)
+create_task_copyfn (gomp_task *task_stmt, omp_context *ctx)
{
struct function *child_cfun;
tree child_fn, t, c, src, dst, f, sf, arg, sarg, decl;
memset (&tcctx, '\0', sizeof (tcctx));
tcctx.cb.src_fn = ctx->cb.src_fn;
tcctx.cb.dst_fn = child_fn;
- tcctx.cb.src_node = cgraph_get_node (tcctx.cb.src_fn);
+ tcctx.cb.src_node = cgraph_node::get (tcctx.cb.src_fn);
gcc_checking_assert (tcctx.cb.src_node);
tcctx.cb.dst_node = tcctx.cb.src_node;
tcctx.cb.src_cfun = ctx->cb.src_cfun;
tcctx.cb.copy_decl = task_copyfn_copy_decl;
tcctx.cb.eh_lp_nr = 0;
tcctx.cb.transform_call_graph_edges = CB_CGE_MOVE;
- tcctx.cb.decl_map = pointer_map_create ();
+ tcctx.cb.decl_map = new hash_map<tree, tree>;
tcctx.ctx = ctx;
if (record_needs_remap)
tree *p;
decl = OMP_CLAUSE_DECL (c);
- p = (tree *) pointer_map_contains (tcctx.cb.decl_map, decl);
+ p = tcctx.cb.decl_map->get (decl);
if (p == NULL)
continue;
n = splay_tree_lookup (ctx->sfield_map, (splay_tree_key) decl);
sf = (tree) n->value;
- sf = *(tree *) pointer_map_contains (tcctx.cb.decl_map, sf);
+ sf = *tcctx.cb.decl_map->get (sf);
src = build_simple_mem_ref_loc (loc, sarg);
src = omp_build_component_ref (src, sf);
t = build2 (MODIFY_EXPR, TREE_TYPE (*p), *p, src);
break;
f = (tree) n->value;
if (tcctx.cb.decl_map)
- f = *(tree *) pointer_map_contains (tcctx.cb.decl_map, f);
+ f = *tcctx.cb.decl_map->get (f);
n = splay_tree_lookup (ctx->sfield_map, (splay_tree_key) decl);
sf = (tree) n->value;
if (tcctx.cb.decl_map)
- sf = *(tree *) pointer_map_contains (tcctx.cb.decl_map, sf);
+ sf = *tcctx.cb.decl_map->get (sf);
src = build_simple_mem_ref_loc (loc, sarg);
src = omp_build_component_ref (src, sf);
dst = build_simple_mem_ref_loc (loc, arg);
break;
f = (tree) n->value;
if (tcctx.cb.decl_map)
- f = *(tree *) pointer_map_contains (tcctx.cb.decl_map, f);
+ f = *tcctx.cb.decl_map->get (f);
n = splay_tree_lookup (ctx->sfield_map, (splay_tree_key) decl);
if (n != NULL)
{
sf = (tree) n->value;
if (tcctx.cb.decl_map)
- sf = *(tree *) pointer_map_contains (tcctx.cb.decl_map, sf);
+ sf = *tcctx.cb.decl_map->get (sf);
src = build_simple_mem_ref_loc (loc, sarg);
src = omp_build_component_ref (src, sf);
if (use_pointer_for_field (decl, NULL) || is_reference (decl))
n = splay_tree_lookup (ctx->field_map, (splay_tree_key) decl);
f = (tree) n->value;
if (tcctx.cb.decl_map)
- f = *(tree *) pointer_map_contains (tcctx.cb.decl_map, f);
+ f = *tcctx.cb.decl_map->get (f);
n = splay_tree_lookup (ctx->sfield_map, (splay_tree_key) decl);
if (n != NULL)
{
sf = (tree) n->value;
if (tcctx.cb.decl_map)
- sf = *(tree *) pointer_map_contains (tcctx.cb.decl_map, sf);
+ sf = *tcctx.cb.decl_map->get (sf);
src = build_simple_mem_ref_loc (loc, sarg);
src = omp_build_component_ref (src, sf);
if (use_pointer_for_field (decl, NULL))
if (n == NULL)
continue;
f = (tree) n->value;
- f = *(tree *) pointer_map_contains (tcctx.cb.decl_map, f);
+ f = *tcctx.cb.decl_map->get (f);
gcc_assert (DECL_HAS_VALUE_EXPR_P (decl));
ind = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (ind) == INDIRECT_REF);
n = splay_tree_lookup (ctx->sfield_map,
(splay_tree_key) TREE_OPERAND (ind, 0));
sf = (tree) n->value;
- sf = *(tree *) pointer_map_contains (tcctx.cb.decl_map, sf);
+ sf = *tcctx.cb.decl_map->get (sf);
src = build_simple_mem_ref_loc (loc, sarg);
src = omp_build_component_ref (src, sf);
src = build_simple_mem_ref_loc (loc, src);
n = splay_tree_lookup (ctx->field_map,
(splay_tree_key) TREE_OPERAND (ind, 0));
df = (tree) n->value;
- df = *(tree *) pointer_map_contains (tcctx.cb.decl_map, df);
+ df = *tcctx.cb.decl_map->get (df);
ptr = build_simple_mem_ref_loc (loc, arg);
ptr = omp_build_component_ref (ptr, df);
t = build2 (MODIFY_EXPR, TREE_TYPE (ptr), ptr,
append_to_statement_list (t, &list);
if (tcctx.cb.decl_map)
- pointer_map_destroy (tcctx.cb.decl_map);
+ delete tcctx.cb.decl_map;
pop_gimplify_context (NULL);
BIND_EXPR_BODY (bind) = list;
pop_cfun ();
gcc_unreachable ();
}
tree type = build_array_type_nelts (ptr_type_node, n_in + n_out + 2);
- tree array = create_tmp_var (type, NULL);
+ tree array = create_tmp_var (type);
tree r = build4 (ARRAY_REF, ptr_type_node, array, size_int (0), NULL_TREE,
NULL_TREE);
g = gimple_build_assign (r, build_int_cst (ptr_type_node, n_in + n_out));
tree clauses;
tree child_fn, t;
gimple stmt = gsi_stmt (*gsi_p);
- gimple par_bind, bind, dep_bind = NULL;
+ gbind *par_bind, *bind, *dep_bind = NULL;
gimple_seq par_body, olist, ilist, par_olist, par_rlist, par_ilist, new_body;
location_t loc = gimple_location (stmt);
clauses = gimple_omp_taskreg_clauses (stmt);
- par_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
+ par_bind
+ = as_a <gbind *> (gimple_seq_first_stmt (gimple_omp_body (stmt)));
par_body = gimple_bind_body (par_bind);
child_fn = ctx->cb.dst_fn;
if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL
}
if (ctx->srecord_type)
- create_task_copyfn (stmt, ctx);
+ create_task_copyfn (as_a <gomp_task *> (stmt), ctx);
push_gimplify_context ();
gimple_seq_add_stmt (&new_body, gimple_build_label (ctx->cancel_label));
gimple_seq_add_seq (&new_body, par_olist);
new_body = maybe_catch_exception (new_body);
+ if (gimple_code (stmt) == GIMPLE_OMP_TASK)
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_omp_continue (integer_zero_node,
+ integer_zero_node));
gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
gimple_omp_set_body (stmt, new_body);
}
}
-/* Lower the OpenMP target directive in the current statement
+/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
static void
{
tree clauses;
tree child_fn, t, c;
- gimple stmt = gsi_stmt (*gsi_p);
- gimple tgt_bind = NULL, bind;
- gimple_seq tgt_body = NULL, olist, ilist, new_body;
+ gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
+ gbind *tgt_bind, *bind;
+ gimple_seq tgt_body, olist, ilist, orlist, irlist, new_body;
location_t loc = gimple_location (stmt);
- int kind = gimple_omp_target_kind (stmt);
+ bool offloaded, data_region;
unsigned int map_cnt = 0;
+ offloaded = is_gimple_omp_offloaded (stmt);
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ data_region = false;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ data_region = true;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
clauses = gimple_omp_target_clauses (stmt);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+
+ tgt_bind = NULL;
+ tgt_body = NULL;
+ if (offloaded)
{
- tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
+ tgt_bind = gimple_seq_first_stmt_as_a_bind (gimple_omp_body (stmt));
tgt_body = gimple_bind_body (tgt_bind);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (data_region)
tgt_body = gimple_omp_body (stmt);
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
+ irlist = NULL;
+ orlist = NULL;
+ if (offloaded
+ && is_gimple_omp_oacc (stmt))
+ oacc_process_reduction_data (&tgt_body, &irlist, &orlist, ctx);
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
default:
break;
case OMP_CLAUSE_MAP:
+#ifdef ENABLE_CHECKING
+ /* First check what we're prepared to handle in the following. */
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO_PSET:
+ break;
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_FORCE_DEALLOC:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ gcc_assert (is_gimple_omp_oacc (stmt));
+ break;
+ default:
+ gcc_unreachable ();
+ }
+#endif
+ /* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
var = OMP_CLAUSE_DECL (c);
if (!maybe_lookup_field (var, ctx))
continue;
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
x = build_simple_mem_ref (x);
map_cnt++;
}
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
target_nesting_level++;
lower_omp (&tgt_body, ctx);
target_nesting_level--;
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (data_region)
lower_omp (&tgt_body, ctx);
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
/* Declare all the variables created by mapping and the variables
declared in the scope of the target body. */
DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
+ tree tkind_type;
+ int talign_shift;
+ if (is_gimple_omp_oacc (stmt))
+ {
+ tkind_type = short_unsigned_type_node;
+ talign_shift = 8;
+ }
+ else
+ {
+ tkind_type = unsigned_char_type_node;
+ talign_shift = 3;
+ }
TREE_VEC_ELT (t, 2)
- = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
- map_cnt),
+ = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
".omp_data_kinds");
DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
continue;
}
+ unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+ if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
+ talign = DECL_ALIGN_UNIT (ovar);
if (nc)
{
tree var = lookup_decl_in_outer_ctx (ovar, ctx);
tree x = build_sender_ref (ovar, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
- && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
- && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ if (maybe_lookup_oacc_reduction (var, ctx))
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
+ gcc_checking_assert (offloaded
+ && is_gimple_omp_oacc (stmt));
+ gimplify_assign (x, var, &ilist);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ {
+ gcc_assert (offloaded);
tree avar
- = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
+ = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
mark_addressable (avar);
gimplify_assign (avar, build_fold_addr_expr (var), &ilist);
+ talign = DECL_ALIGN_UNIT (avar);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
}
else if (is_gimple_reg (var))
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
- tree avar = create_tmp_var (TREE_TYPE (var), NULL);
+ gcc_assert (offloaded);
+ tree avar = create_tmp_var (TREE_TYPE (var));
mark_addressable (avar);
- if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
- && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+ enum gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (c);
+ if (GOMP_MAP_COPY_TO_P (map_kind)
+ || map_kind == GOMP_MAP_POINTER
+ || map_kind == GOMP_MAP_TO_PSET
+ || map_kind == GOMP_MAP_FORCE_DEVICEPTR)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
- || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+ if ((GOMP_MAP_COPY_FROM_P (map_kind)
+ || map_kind == GOMP_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
- unsigned char tkind = 0;
+ unsigned HOST_WIDE_INT tkind;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_MAP:
tkind = OMP_CLAUSE_MAP_KIND (c);
break;
case OMP_CLAUSE_TO:
- tkind = OMP_CLAUSE_MAP_TO;
+ tkind = GOMP_MAP_TO;
break;
case OMP_CLAUSE_FROM:
- tkind = OMP_CLAUSE_MAP_FROM;
+ tkind = GOMP_MAP_FROM;
break;
default:
gcc_unreachable ();
}
- unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
- if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
- talign = DECL_ALIGN_UNIT (ovar);
+ gcc_checking_assert (tkind
+ < (HOST_WIDE_INT_C (1U) << talign_shift));
talign = ceil_log2 (talign);
- tkind |= talign << 3;
+ tkind |= talign << talign_shift;
+ gcc_checking_assert (tkind
+ <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
- build_int_cst (unsigned_char_type_node,
- tkind));
+ build_int_cstu (tkind_type, tkind));
if (nc && nc != c)
c = nc;
}
new_body = NULL;
- if (ctx->record_type && kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded
+ && ctx->record_type)
{
t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
/* fixup_child_record_type might have changed receiver_decl's type. */
gimple_build_assign (ctx->receiver_decl, t));
}
- if (kind == GF_OMP_TARGET_KIND_REGION)
+ if (offloaded)
{
gimple_seq_add_seq (&new_body, tgt_body);
new_body = maybe_catch_exception (new_body);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (data_region)
new_body = tgt_body;
- if (kind != GF_OMP_TARGET_KIND_UPDATE)
+ if (offloaded || data_region)
{
gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
gimple_omp_set_body (stmt, new_body);
tgt_bind ? gimple_bind_block (tgt_bind)
: NULL_TREE);
gsi_replace (gsi_p, bind, true);
+ gimple_bind_add_seq (bind, irlist);
gimple_bind_add_seq (bind, ilist);
gimple_bind_add_stmt (bind, stmt);
gimple_bind_add_seq (bind, olist);
+ gimple_bind_add_seq (bind, orlist);
pop_gimplify_context (NULL);
}
static void
lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
- gimple teams_stmt = gsi_stmt (*gsi_p);
+ gomp_teams *teams_stmt = as_a <gomp_teams *> (gsi_stmt (*gsi_p));
push_gimplify_context ();
tree block = make_node (BLOCK);
- gimple bind = gimple_build_bind (NULL, NULL, block);
+ gbind *bind = gimple_build_bind (NULL, NULL, block);
gsi_replace (gsi_p, bind, true);
gimple_seq bind_body = NULL;
gimple_seq dlist = NULL;
/* Callback for lower_omp_1. Return non-NULL if *tp needs to be
regimplified. If DATA is non-NULL, lower_omp_1 is outside
- of OpenMP context, but with task_shared_vars set. */
+ of OMP context, but with task_shared_vars set. */
static tree
lower_omp_regimplify_p (tree *tp, int *walk_subtrees,
{
gimple stmt = gsi_stmt (*gsi_p);
struct walk_stmt_info wi;
+ gcall *call_stmt;
if (gimple_has_location (stmt))
input_location = gimple_location (stmt);
memset (&wi, '\0', sizeof (wi));
/* If we have issued syntax errors, avoid doing any heavy lifting.
- Just replace the OpenMP directives with a NOP to avoid
+ Just replace the OMP directives with a NOP to avoid
confusing RTL expansion. */
if (seen_error () && is_gimple_omp (stmt))
{
switch (gimple_code (stmt))
{
case GIMPLE_COND:
- if ((ctx || task_shared_vars)
- && (walk_tree (gimple_cond_lhs_ptr (stmt), lower_omp_regimplify_p,
- ctx ? NULL : &wi, NULL)
- || walk_tree (gimple_cond_rhs_ptr (stmt), lower_omp_regimplify_p,
- ctx ? NULL : &wi, NULL)))
- gimple_regimplify_operands (stmt, gsi_p);
+ {
+ gcond *cond_stmt = as_a <gcond *> (stmt);
+ if ((ctx || task_shared_vars)
+ && (walk_tree (gimple_cond_lhs_ptr (cond_stmt),
+ lower_omp_regimplify_p,
+ ctx ? NULL : &wi, NULL)
+ || walk_tree (gimple_cond_rhs_ptr (cond_stmt),
+ lower_omp_regimplify_p,
+ ctx ? NULL : &wi, NULL)))
+ gimple_regimplify_operands (cond_stmt, gsi_p);
+ }
break;
case GIMPLE_CATCH:
- lower_omp (gimple_catch_handler_ptr (stmt), ctx);
+ lower_omp (gimple_catch_handler_ptr (as_a <gcatch *> (stmt)), ctx);
break;
case GIMPLE_EH_FILTER:
lower_omp (gimple_eh_filter_failure_ptr (stmt), ctx);
lower_omp (gimple_try_cleanup_ptr (stmt), ctx);
break;
case GIMPLE_TRANSACTION:
- lower_omp (gimple_transaction_body_ptr (stmt), ctx);
+ lower_omp (gimple_transaction_body_ptr (
+ as_a <gtransaction *> (stmt)),
+ ctx);
break;
case GIMPLE_BIND:
- lower_omp (gimple_bind_body_ptr (stmt), ctx);
+ lower_omp (gimple_bind_body_ptr (as_a <gbind *> (stmt)), ctx);
break;
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
break;
case GIMPLE_OMP_ATOMIC_LOAD:
if ((ctx || task_shared_vars)
- && walk_tree (gimple_omp_atomic_load_rhs_ptr (stmt),
+ && walk_tree (gimple_omp_atomic_load_rhs_ptr (
+ as_a <gomp_atomic_load *> (stmt)),
lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
gimple_regimplify_operands (stmt, gsi_p);
break;
break;
case GIMPLE_CALL:
tree fndecl;
- fndecl = gimple_call_fndecl (stmt);
+ call_stmt = as_a <gcall *> (stmt);
+ fndecl = gimple_call_fndecl (call_stmt);
if (fndecl
&& DECL_BUILT_IN_CLASS (fndecl) == BUILT_IN_NORMAL)
switch (DECL_FUNCTION_CODE (fndecl))
cctx = ctx;
if (gimple_code (cctx->stmt) == GIMPLE_OMP_SECTION)
cctx = cctx->outer;
- gcc_assert (gimple_call_lhs (stmt) == NULL_TREE);
+ gcc_assert (gimple_call_lhs (call_stmt) == NULL_TREE);
if (!cctx->cancellable)
{
if (DECL_FUNCTION_CODE (fndecl)
if (DECL_FUNCTION_CODE (fndecl) == BUILT_IN_GOMP_BARRIER)
{
fndecl = builtin_decl_explicit (BUILT_IN_GOMP_BARRIER_CANCEL);
- gimple_call_set_fndecl (stmt, fndecl);
- gimple_call_set_fntype (stmt, TREE_TYPE (fndecl));
+ gimple_call_set_fndecl (call_stmt, fndecl);
+ gimple_call_set_fntype (call_stmt, TREE_TYPE (fndecl));
}
tree lhs;
- lhs = create_tmp_var (TREE_TYPE (TREE_TYPE (fndecl)), NULL);
- gimple_call_set_lhs (stmt, lhs);
+ lhs = create_tmp_var (TREE_TYPE (TREE_TYPE (fndecl)));
+ gimple_call_set_lhs (call_stmt, lhs);
tree fallthru_label;
fallthru_label = create_artificial_label (UNKNOWN_LOCATION);
gimple g;
gimple_stmt_iterator gsi;
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
lower_omp_1 (&gsi, ctx);
- /* During gimplification, we have not always invoked fold_stmt
- (gimplify.c:maybe_fold_stmt); call it now. */
+ /* During gimplification, we haven't folded statments inside offloading
+ regions (gimplify.c:maybe_fold_stmt); do that now. */
if (target_nesting_level)
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
fold_stmt (&gsi);
execute_lower_omp (void)
{
gimple_seq body;
+ int i;
+ omp_context *ctx;
/* This pass always runs, to provide PROP_gimple_lomp.
- But there is nothing to do unless -fopenmp is given. */
- if (flag_openmp == 0 && flag_openmp_simd == 0 && flag_cilkplus == 0)
+ But often, there is nothing to do. */
+ if (flag_cilkplus == 0 && flag_openacc == 0 && flag_openmp == 0
+ && flag_openmp_simd == 0)
return 0;
all_contexts = splay_tree_new (splay_tree_compare_pointers, 0,
body = gimple_body (current_function_decl);
scan_omp (&body, NULL);
gcc_assert (taskreg_nesting_level == 0);
+ FOR_EACH_VEC_ELT (taskreg_contexts, i, ctx)
+ finish_taskreg_scan (ctx);
+ taskreg_contexts.release ();
if (all_contexts->root)
{
GIMPLE_PASS, /* type */
"omplower", /* name */
OPTGROUP_NONE, /* optinfo_flags */
- true, /* has_execute */
TV_NONE, /* tv_id */
PROP_gimple_any, /* properties_required */
PROP_gimple_lomp, /* properties_provided */
return new pass_lower_omp (ctxt);
}
\f
-/* The following is a utility to diagnose OpenMP structured block violations.
+/* The following is a utility to diagnose structured block violations.
It is not part of the "omplower" pass, as that's invoked too late. It
should be invoked by the respective front ends after gimplification. */
diagnose_sb_0 (gimple_stmt_iterator *gsi_p,
gimple branch_ctx, gimple label_ctx)
{
+ gcc_checking_assert (!branch_ctx || is_gimple_omp (branch_ctx));
+ gcc_checking_assert (!label_ctx || is_gimple_omp (label_ctx));
+
if (label_ctx == branch_ctx)
return false;
+ const char* kind = NULL;
+
+ if (flag_cilkplus)
+ {
+ if ((branch_ctx
+ && gimple_code (branch_ctx) == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (branch_ctx) == GF_OMP_FOR_KIND_CILKSIMD)
+ || (label_ctx
+ && gimple_code (label_ctx) == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (label_ctx) == GF_OMP_FOR_KIND_CILKSIMD))
+ kind = "Cilk Plus";
+ }
+ if (flag_openacc)
+ {
+ if ((branch_ctx && is_gimple_omp_oacc (branch_ctx))
+ || (label_ctx && is_gimple_omp_oacc (label_ctx)))
+ {
+ gcc_checking_assert (kind == NULL);
+ kind = "OpenACC";
+ }
+ }
+ if (kind == NULL)
+ {
+ gcc_checking_assert (flag_openmp);
+ kind = "OpenMP";
+ }
/*
Previously we kept track of the label's entire context in diagnose_sb_[12]
}
if (exit_p)
- error ("invalid exit from OpenMP structured block");
+ error ("invalid exit from %s structured block", kind);
else
- error ("invalid entry to OpenMP structured block");
+ error ("invalid entry to %s structured block", kind);
#endif
- bool cilkplus_block = false;
- if (flag_cilkplus)
- {
- if ((branch_ctx
- && gimple_code (branch_ctx) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (branch_ctx) == GF_OMP_FOR_KIND_CILKSIMD)
- || (label_ctx
- && gimple_code (label_ctx) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (label_ctx) == GF_OMP_FOR_KIND_CILKSIMD))
- cilkplus_block = true;
- }
-
/* If it's obvious we have an invalid entry, be specific about the error. */
if (branch_ctx == NULL)
- {
- if (cilkplus_block)
- error ("invalid entry to Cilk Plus structured block");
- else
- error ("invalid entry to OpenMP structured block");
- }
+ error ("invalid entry to %s structured block", kind);
else
{
/* Otherwise, be vague and lazy, but efficient. */
- if (cilkplus_block)
- error ("invalid branch to/from a Cilk Plus structured block");
- else
- error ("invalid branch to/from an OpenMP structured block");
+ error ("invalid branch to/from %s structured block", kind);
}
gsi_replace (gsi_p, gimple_build_nop (), false);
return true;
}
-/* Pass 1: Create a minimal tree of OpenMP structured blocks, and record
+/* Pass 1: Create a minimal tree of structured blocks, and record
where each label is found. */
static tree
*handled_ops_p = true;
- switch (gimple_code (stmt))
+ switch (gimple_code (stmt))
{
WALK_SUBSTMTS;
break;
case GIMPLE_LABEL:
- splay_tree_insert (all_labels, (splay_tree_key) gimple_label_label (stmt),
+ splay_tree_insert (all_labels,
+ (splay_tree_key) gimple_label_label (
+ as_a <glabel *> (stmt)),
(splay_tree_value) context);
break;
case GIMPLE_COND:
{
- tree lab = gimple_cond_true_label (stmt);
+ gcond *cond_stmt = as_a <gcond *> (stmt);
+ tree lab = gimple_cond_true_label (cond_stmt);
if (lab)
{
n = splay_tree_lookup (all_labels,
diagnose_sb_0 (gsi_p, context,
n ? (gimple) n->value : NULL);
}
- lab = gimple_cond_false_label (stmt);
+ lab = gimple_cond_false_label (cond_stmt);
if (lab)
{
n = splay_tree_lookup (all_labels,
case GIMPLE_SWITCH:
{
+ gswitch *switch_stmt = as_a <gswitch *> (stmt);
unsigned int i;
- for (i = 0; i < gimple_switch_num_labels (stmt); ++i)
+ for (i = 0; i < gimple_switch_num_labels (switch_stmt); ++i)
{
- tree lab = CASE_LABEL (gimple_switch_label (stmt, i));
+ tree lab = CASE_LABEL (gimple_switch_label (switch_stmt, i));
n = splay_tree_lookup (all_labels, (splay_tree_key) lab);
if (n && diagnose_sb_0 (gsi_p, context, (gimple) n->value))
break;
return NULL_TREE;
}
-/* Called from tree-cfg.c::make_edges to create cfg edges for all GIMPLE_OMP
- codes. */
+/* Called from tree-cfg.c::make_edges to create cfg edges for all relevant
+ GIMPLE_* codes. */
bool
make_gimple_omp_edges (basic_block bb, struct omp_region **region,
int *region_idx)
case GIMPLE_OMP_TARGET:
cur_region = new_omp_region (bb, code, cur_region);
fallthru = true;
- if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE)
- cur_region = cur_region->outer;
+ switch (gimple_omp_target_kind (last))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ cur_region = cur_region->outer;
+ break;
+ default:
+ gcc_unreachable ();
+ }
break;
case GIMPLE_OMP_SECTIONS:
somewhere other than the next block. This will be
created later. */
cur_region->exit = bb;
+ if (cur_region->type == GIMPLE_OMP_TASK)
+ /* Add an edge corresponding to not scheduling the task
+ immediately. */
+ make_edge (cur_region->entry, bb, EDGE_ABNORMAL);
fallthru = cur_region->type != GIMPLE_OMP_SECTION;
cur_region = cur_region->outer;
break;
}
break;
+ case GIMPLE_OMP_TASK:
+ fallthru = true;
+ break;
+
default:
gcc_unreachable ();
}
GIMPLE_PASS, /* type */
"*diagnose_omp_blocks", /* name */
OPTGROUP_NONE, /* optinfo_flags */
- true, /* has_execute */
TV_NONE, /* tv_id */
PROP_gimple_any, /* properties_required */
0, /* properties_provided */
{}
/* opt_pass methods: */
- virtual bool gate (function *) { return flag_openmp || flag_cilkplus; }
+ virtual bool gate (function *)
+ {
+ return flag_cilkplus || flag_openacc || flag_openmp;
+ }
virtual unsigned int execute (function *)
{
return diagnose_omp_structured_block_errors ();
}
pp_underscore (&pp);
- pp_string (&pp,
- IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (node->decl)));
- const char *str = pp_formatted_text (&pp);
+ const char *str = IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (node->decl));
+ if (*str == '*')
+ ++str;
+ pp_string (&pp, str);
+ str = pp_formatted_text (&pp);
/* If there already is a SIMD clone with the same mangled name, don't
add another one. This can happen e.g. for
struct cgraph_node *new_node;
if (old_node->definition)
{
- if (!cgraph_function_with_gimple_body_p (old_node))
+ if (!old_node->has_gimple_body_p ())
return NULL;
- cgraph_get_body (old_node);
- new_node = cgraph_function_versioning (old_node, vNULL, NULL, NULL,
- false, NULL, NULL, "simdclone");
+ old_node->get_body ();
+ new_node = old_node->create_version_clone_with_body (vNULL, NULL, NULL,
+ false, NULL, NULL,
+ "simdclone");
}
else
{
SET_DECL_RTL (new_decl, NULL);
DECL_STATIC_CONSTRUCTOR (new_decl) = 0;
DECL_STATIC_DESTRUCTOR (new_decl) = 0;
- new_node
- = cgraph_copy_node_for_versioning (old_node, new_decl, vNULL, NULL);
- cgraph_call_function_insertion_hooks (new_node);
+ new_node = old_node->create_version_clone (new_decl, vNULL, NULL);
+ symtab->call_cgraph_insertion_hooks (new_node);
}
if (new_node == NULL)
return new_node;
if (orig_rettype == void_type_node)
return NULL_TREE;
TREE_TYPE (fndecl) = build_distinct_type_copy (TREE_TYPE (fndecl));
- if (INTEGRAL_TYPE_P (TREE_TYPE (TREE_TYPE (fndecl)))
- || POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (fndecl))))
+ t = TREE_TYPE (TREE_TYPE (fndecl));
+ if (INTEGRAL_TYPE_P (t) || POINTER_TYPE_P (t))
veclen = node->simdclone->vecsize_int;
else
veclen = node->simdclone->vecsize_float;
- veclen /= GET_MODE_BITSIZE (TYPE_MODE (TREE_TYPE (TREE_TYPE (fndecl))));
+ veclen /= GET_MODE_BITSIZE (TYPE_MODE (t));
if (veclen > node->simdclone->simdlen)
veclen = node->simdclone->simdlen;
+ if (POINTER_TYPE_P (t))
+ t = pointer_sized_int_node;
if (veclen == node->simdclone->simdlen)
- TREE_TYPE (TREE_TYPE (fndecl))
- = build_vector_type (TREE_TYPE (TREE_TYPE (fndecl)),
- node->simdclone->simdlen);
+ t = build_vector_type (t, node->simdclone->simdlen);
else
{
- t = build_vector_type (TREE_TYPE (TREE_TYPE (fndecl)), veclen);
+ t = build_vector_type (t, veclen);
t = build_array_type_nelts (t, node->simdclone->simdlen / veclen);
- TREE_TYPE (TREE_TYPE (fndecl)) = t;
}
+ TREE_TYPE (TREE_TYPE (fndecl)) = t;
if (!node->definition)
return NULL_TREE;
if (veclen > node->simdclone->simdlen)
veclen = node->simdclone->simdlen;
adj.arg_prefix = "simd";
- adj.type = build_vector_type (parm_type, veclen);
+ if (POINTER_TYPE_P (parm_type))
+ adj.type = build_vector_type (pointer_sized_int_node, veclen);
+ else
+ adj.type = build_vector_type (parm_type, veclen);
node->simdclone->args[i].vector_type = adj.type;
for (j = veclen; j < node->simdclone->simdlen; j += veclen)
{
veclen /= GET_MODE_BITSIZE (TYPE_MODE (base_type));
if (veclen > node->simdclone->simdlen)
veclen = node->simdclone->simdlen;
- adj.type = build_vector_type (base_type, veclen);
+ if (POINTER_TYPE_P (base_type))
+ adj.type = build_vector_type (pointer_sized_int_node, veclen);
+ else
+ adj.type = build_vector_type (base_type, veclen);
adjustments.safe_push (adj);
for (j = veclen; j < node->simdclone->simdlen; j += veclen)
if (tp != orig_tp)
{
repl = build_fold_addr_expr (repl);
- gimple stmt
- = gimple_build_assign (make_ssa_name (TREE_TYPE (repl), NULL), repl);
- repl = gimple_assign_lhs (stmt);
+ gimple stmt;
+ if (is_gimple_debug (info->stmt))
+ {
+ tree vexpr = make_node (DEBUG_EXPR_DECL);
+ stmt = gimple_build_debug_source_bind (vexpr, repl, NULL);
+ DECL_ARTIFICIAL (vexpr) = 1;
+ TREE_TYPE (vexpr) = TREE_TYPE (repl);
+ DECL_MODE (vexpr) = TYPE_MODE (TREE_TYPE (repl));
+ repl = vexpr;
+ }
+ else
+ {
+ stmt = gimple_build_assign (make_ssa_name (TREE_TYPE (repl)), repl);
+ repl = gimple_assign_lhs (stmt);
+ }
gimple_stmt_iterator gsi = gsi_for_stmt (info->stmt);
gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
*orig_tp = repl;
wi.info = &info;
walk_gimple_op (stmt, ipa_simd_modify_stmt_ops, &wi);
- if (gimple_code (stmt) == GIMPLE_RETURN)
+ if (greturn *return_stmt = dyn_cast <greturn *> (stmt))
{
- tree retval = gimple_return_retval (stmt);
+ tree retval = gimple_return_retval (return_stmt);
if (!retval)
{
gsi_remove (&gsi, true);
/* Adjust all uses of vector arguments accordingly. Adjust all
return values accordingly. */
tree iter = create_tmp_var (unsigned_type_node, "iter");
- tree iter1 = make_ssa_name (iter, NULL);
- tree iter2 = make_ssa_name (iter, NULL);
+ tree iter1 = make_ssa_name (iter);
+ tree iter2 = make_ssa_name (iter);
ipa_simd_modify_function_body (node, adjustments, retval, iter1);
/* Initialize the iteration variable. */
iteration increment and the condition/branch. */
basic_block orig_exit = EDGE_PRED (EXIT_BLOCK_PTR_FOR_FN (cfun), 0)->src;
basic_block incr_bb = create_empty_bb (orig_exit);
+ add_bb_to_loop (incr_bb, body_bb->loop_father);
/* The succ of orig_exit was EXIT_BLOCK_PTR_FOR_FN (cfun), with an empty
flag. Set it now to be a FALLTHRU_EDGE. */
gcc_assert (EDGE_COUNT (orig_exit->succs) == 1);
edge e = make_edge (incr_bb, EXIT_BLOCK_PTR_FOR_FN (cfun), 0);
e->probability = REG_BR_PROB_BASE;
gsi = gsi_last_bb (incr_bb);
- gimple g = gimple_build_assign_with_ops (PLUS_EXPR, iter2, iter1,
- build_int_cst (unsigned_type_node,
- 1));
+ gimple g = gimple_build_assign (iter2, PLUS_EXPR, iter1,
+ build_int_cst (unsigned_type_node, 1));
gsi_insert_after (&gsi, g, GSI_CONTINUE_LINKING);
/* Mostly annotate the loop for the vectorizer (the rest is done below). */
loop->safelen = node->simdclone->simdlen;
loop->force_vectorize = true;
loop->header = body_bb;
- add_bb_to_loop (incr_bb, loop);
/* Branch around the body if the mask applies. */
if (node->simdclone->inbranch)
gimple_stmt_iterator gsi = gsi_last_bb (loop->header);
tree mask_array
= node->simdclone->args[node->simdclone->nargs - 1].simd_array;
- tree mask = make_ssa_name (TREE_TYPE (TREE_TYPE (mask_array)), NULL);
+ tree mask = make_ssa_name (TREE_TYPE (TREE_TYPE (mask_array)));
tree aref = build4 (ARRAY_REF,
TREE_TYPE (TREE_TYPE (mask_array)),
mask_array, iter1,
{
aref = build1 (VIEW_CONVERT_EXPR,
build_nonstandard_integer_type (bitsize, 0), mask);
- mask = make_ssa_name (TREE_TYPE (aref), NULL);
+ mask = make_ssa_name (TREE_TYPE (aref));
g = gimple_build_assign (mask, aref);
gsi_insert_after (&gsi, g, GSI_CONTINUE_LINKING);
}
gsi_insert_after (&gsi, g, GSI_CONTINUE_LINKING);
e = split_block (incr_bb, gsi_stmt (gsi));
basic_block latch_bb = e->dest;
- basic_block new_exit_bb = e->dest;
- new_exit_bb = split_block (latch_bb, NULL)->dest;
+ basic_block new_exit_bb;
+ new_exit_bb = split_block_after_labels (latch_bb)->dest;
loop->latch = latch_bb;
redirect_edge_succ (FALLTHRU_EDGE (latch_bb), body_bb);
make_edge (incr_bb, latch_bb, EDGE_TRUE_VALUE); */
FALLTHRU_EDGE (incr_bb)->flags = EDGE_TRUE_VALUE;
- gimple phi = create_phi_node (iter1, body_bb);
+ gphi *phi = create_phi_node (iter1, body_bb);
edge preheader_edge = find_edge (entry_bb, body_bb);
edge latch_edge = single_succ_edge (latch_bb);
add_phi_arg (phi, build_zero_cst (unsigned_type_node), preheader_edge,
tree fn = builtin_decl_explicit (BUILT_IN_ASSUME_ALIGNED);
gimple_seq seq = NULL;
bool need_cvt = false;
- gimple call
+ gcall *call
= gimple_build_call (fn, 2, def, size_int (alignment));
g = call;
if (!useless_type_conversion_p (TREE_TYPE (orig_arg),
ptr_type_node))
need_cvt = true;
- tree t = make_ssa_name (need_cvt ? ptr_type_node : orig_arg, NULL);
+ tree t = make_ssa_name (need_cvt ? ptr_type_node : orig_arg);
gimple_call_set_lhs (g, t);
gimple_seq_add_stmt_without_update (&seq, g);
if (need_cvt)
{
- t = make_ssa_name (orig_arg, NULL);
- g = gimple_build_assign_with_ops (NOP_EXPR, t,
- gimple_call_lhs (g),
- NULL_TREE);
+ t = make_ssa_name (orig_arg);
+ g = gimple_build_assign (t, NOP_EXPR, gimple_call_lhs (g));
gimple_seq_add_stmt_without_update (&seq, g);
}
gsi_insert_seq_on_edge_immediate
entry_bb = single_succ (ENTRY_BLOCK_PTR_FOR_FN (cfun));
int freq = compute_call_stmt_bb_frequency (current_function_decl,
entry_bb);
- cgraph_create_edge (node, cgraph_get_create_node (fn),
- call, entry_bb->count, freq);
+ node->create_edge (cgraph_node::get_create (fn),
+ call, entry_bb->count, freq);
imm_use_iterator iter;
use_operand_p use_p;
|| POINTER_TYPE_P (TREE_TYPE (orig_arg)));
if (def && !has_zero_uses (def))
{
- iter1 = make_ssa_name (orig_arg, NULL);
- iter2 = make_ssa_name (orig_arg, NULL);
+ iter1 = make_ssa_name (orig_arg);
+ iter2 = make_ssa_name (orig_arg);
phi = create_phi_node (iter1, body_bb);
add_phi_arg (phi, def, preheader_edge, UNKNOWN_LOCATION);
add_phi_arg (phi, iter2, latch_edge, UNKNOWN_LOCATION);
? TREE_TYPE (orig_arg) : sizetype;
tree addcst
= build_int_cst (addtype, node->simdclone->args[i].linear_step);
- g = gimple_build_assign_with_ops (code, iter2, iter1, addcst);
+ g = gimple_build_assign (iter2, code, iter1, addcst);
gsi = gsi_last_bb (incr_bb);
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
clone->prev_clone->simdclone->next_clone = n;
node->simd_clones->simdclone->prev_clone = n;
}
- change_decl_assembler_name (n->decl, id);
+ symtab->change_decl_assembler_name (n->decl, id);
/* And finally adjust the return type, parameters and for
definitions also function body. */
if (node->definition)
SIMPLE_IPA_PASS, /* type */
"simdclone", /* name */
OPTGROUP_NONE, /* optinfo_flags */
- true, /* has_execute */
TV_NONE, /* tv_id */
( PROP_ssa | PROP_cfg ), /* properties_required */
0, /* properties_provided */
return new pass_omp_simd_clone (ctxt);
}
+/* Helper function for omp_finish_file routine. Takes decls from V_DECLS and
+ adds their addresses and sizes to constructor-vector V_CTOR. */
+static void
+add_decls_addresses_to_decl_constructor (vec<tree, va_gc> *v_decls,
+ vec<constructor_elt, va_gc> *v_ctor)
+{
+ unsigned len = vec_safe_length (v_decls);
+ for (unsigned i = 0; i < len; i++)
+ {
+ tree it = (*v_decls)[i];
+ bool is_function = TREE_CODE (it) != VAR_DECL;
+
+ CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE, build_fold_addr_expr (it));
+ if (!is_function)
+ CONSTRUCTOR_APPEND_ELT (v_ctor, NULL_TREE,
+ fold_convert (const_ptr_type_node,
+ DECL_SIZE_UNIT (it)));
+ }
+}
+
+/* Create new symbols containing (address, size) pairs for global variables,
+ marked with "omp declare target" attribute, as well as addresses for the
+ functions, which are outlined offloading regions. */
+void
+omp_finish_file (void)
+{
+ unsigned num_funcs = vec_safe_length (offload_funcs);
+ unsigned num_vars = vec_safe_length (offload_vars);
+
+ if (num_funcs == 0 && num_vars == 0)
+ return;
+
+ if (targetm_common.have_named_sections)
+ {
+ vec<constructor_elt, va_gc> *v_f, *v_v;
+ vec_alloc (v_f, num_funcs);
+ vec_alloc (v_v, num_vars * 2);
+
+ add_decls_addresses_to_decl_constructor (offload_funcs, v_f);
+ add_decls_addresses_to_decl_constructor (offload_vars, v_v);
+
+ tree vars_decl_type = build_array_type_nelts (pointer_sized_int_node,
+ num_vars * 2);
+ tree funcs_decl_type = build_array_type_nelts (pointer_sized_int_node,
+ num_funcs);
+ TYPE_ALIGN (vars_decl_type) = TYPE_ALIGN (pointer_sized_int_node);
+ TYPE_ALIGN (funcs_decl_type) = TYPE_ALIGN (pointer_sized_int_node);
+ tree ctor_v = build_constructor (vars_decl_type, v_v);
+ tree ctor_f = build_constructor (funcs_decl_type, v_f);
+ TREE_CONSTANT (ctor_v) = TREE_CONSTANT (ctor_f) = 1;
+ TREE_STATIC (ctor_v) = TREE_STATIC (ctor_f) = 1;
+ tree funcs_decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+ get_identifier (".offload_func_table"),
+ funcs_decl_type);
+ tree vars_decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+ get_identifier (".offload_var_table"),
+ vars_decl_type);
+ TREE_STATIC (funcs_decl) = TREE_STATIC (vars_decl) = 1;
+ /* Do not align tables more than TYPE_ALIGN (pointer_sized_int_node),
+ otherwise a joint table in a binary will contain padding between
+ tables from multiple object files. */
+ DECL_USER_ALIGN (funcs_decl) = DECL_USER_ALIGN (vars_decl) = 1;
+ DECL_ALIGN (funcs_decl) = TYPE_ALIGN (funcs_decl_type);
+ DECL_ALIGN (vars_decl) = TYPE_ALIGN (vars_decl_type);
+ DECL_INITIAL (funcs_decl) = ctor_f;
+ DECL_INITIAL (vars_decl) = ctor_v;
+ set_decl_section_name (funcs_decl, OFFLOAD_FUNC_TABLE_SECTION_NAME);
+ set_decl_section_name (vars_decl, OFFLOAD_VAR_TABLE_SECTION_NAME);
+
+ varpool_node::finalize_decl (vars_decl);
+ varpool_node::finalize_decl (funcs_decl);
+ }
+ else
+ {
+ for (unsigned i = 0; i < num_funcs; i++)
+ {
+ tree it = (*offload_funcs)[i];
+ targetm.record_offload_symbol (it);
+ }
+ for (unsigned i = 0; i < num_vars; i++)
+ {
+ tree it = (*offload_vars)[i];
+ targetm.record_offload_symbol (it);
+ }
+ }
+}
+
#include "gt-omp-low.h"
projects / gcc.git / blobdiff