Contributed by Diego Novillo <dnovillo@redhat.com>
- Copyright (C) 2005-2019 Free Software Foundation, Inc.
+ Copyright (C) 2005-2020 Free Software Foundation, Inc.
This file is part of GCC.
#include "splay-tree.h"
#include "omp-general.h"
#include "omp-low.h"
-#include "omp-grid.h"
#include "gimple-low.h"
+#include "alloc-pool.h"
#include "symbol-summary.h"
#include "tree-nested.h"
#include "context.h"
#include "gomp-constants.h"
#include "gimple-pretty-print.h"
-#include "hsa-common.h"
#include "stringpool.h"
#include "attribs.h"
corresponding tracking loop iteration variables. */
hash_map<tree, tree> *lastprivate_conditional_map;
+ /* And a hash map from the allocate variables to their corresponding
+ allocators. */
+ hash_map<tree, tree> *allocate_map;
+
+ /* A tree_list of the reduction clauses in this context. This is
+ only used for checking the consistency of OpenACC reduction
+ clauses in scan_omp_for and is not guaranteed to contain a valid
+ value outside of this function. */
+ tree local_reduction_clauses;
+
+ /* A tree_list of the reduction clauses in outer contexts. This is
+ only used for checking the consistency of OpenACC reduction
+ clauses in scan_omp_for and is not guaranteed to contain a valid
+ value outside of this function. */
+ tree outer_reduction_clauses;
+
/* Nesting depth of this context. Used to beautify error messages re
invalid gotos. The outermost ctx is depth 1, with depth 0 being
reserved for the main body of the function. */
/* True in the second simd loop of for simd with inscan reductions. */
bool for_simd_scan_phase;
+
+ /* True if there is order(concurrent) clause on the construct. */
+ bool order_concurrent;
+
+ /* True if there is bind clause on the construct (i.e. a loop construct). */
+ bool loop_p;
};
static splay_tree all_contexts;
static int taskreg_nesting_level;
static int target_nesting_level;
static bitmap task_shared_vars;
+static bitmap global_nonaddressable_vars;
static vec<omp_context *> taskreg_contexts;
static void scan_omp (gimple_seq *, omp_context *);
*handled_ops_p = false; \
break;
-/* Return true if CTX corresponds to an oacc parallel region. */
+/* Return whether CTX represents an OpenACC 'parallel' or 'serial' construct.
+ (This doesn't include OpenACC 'kernels' decomposed parts.) */
static bool
-is_oacc_parallel (omp_context *ctx)
+is_oacc_parallel_or_serial (omp_context *ctx)
{
enum gimple_code outer_type = gimple_code (ctx->stmt);
return ((outer_type == GIMPLE_OMP_TARGET)
- && (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_OACC_PARALLEL));
+ && ((gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ || (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_SERIAL)));
}
-/* Return true if CTX corresponds to an oacc kernels region. */
+/* Return whether CTX represents an OpenACC 'kernels' construct.
+ (This doesn't include OpenACC 'kernels' decomposed parts.) */
static bool
is_oacc_kernels (omp_context *ctx)
== GF_OMP_TARGET_KIND_OACC_KERNELS));
}
+/* Return whether CTX represents an OpenACC 'kernels' decomposed part. */
+
+static bool
+is_oacc_kernels_decomposed_part (omp_context *ctx)
+{
+ enum gimple_code outer_type = gimple_code (ctx->stmt);
+ return ((outer_type == GIMPLE_OMP_TARGET)
+ && ((gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED)
+ || (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE)
+ || (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_DATA_KERNELS)));
+}
+
+/* Return true if STMT corresponds to an OpenMP target region. */
+static bool
+is_omp_target (gimple *stmt)
+{
+ if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+ {
+ int kind = gimple_omp_target_kind (stmt);
+ return (kind == GF_OMP_TARGET_KIND_REGION
+ || kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_ENTER_DATA
+ || kind == GF_OMP_TARGET_KIND_EXIT_DATA);
+ }
+ return false;
+}
+
/* If DECL is the artificial dummy VAR_DECL created for non-static
data member privatization, return the underlying "this" parameter,
otherwise return NULL. */
/* Do not use copy-in/copy-out for variables that have their
address taken. */
- if (TREE_ADDRESSABLE (decl))
+ if (is_global_var (decl))
+ {
+ /* For file scope vars, track whether we've seen them as
+ non-addressable initially and in that case, keep the same
+ answer for the duration of the pass, even when they are made
+ addressable later on e.g. through reduction expansion. Global
+ variables which weren't addressable before the pass will not
+ have their privatized copies address taken. See PR91216. */
+ if (!TREE_ADDRESSABLE (decl))
+ {
+ if (!global_nonaddressable_vars)
+ global_nonaddressable_vars = BITMAP_ALLOC (NULL);
+ bitmap_set_bit (global_nonaddressable_vars, DECL_UID (decl));
+ }
+ else if (!global_nonaddressable_vars
+ || !bitmap_bit_p (global_nonaddressable_vars,
+ DECL_UID (decl)))
+ return true;
+ }
+ else if (TREE_ADDRESSABLE (decl))
return true;
/* lower_send_shared_vars only uses copy-in, but not copy-out
omp_context *up;
for (up = shared_ctx->outer; up; up = up->outer)
- if (is_taskreg_ctx (up) && maybe_lookup_decl (decl, up))
+ if ((is_taskreg_ctx (up)
+ || (gimple_code (up->stmt) == GIMPLE_OMP_TARGET
+ && is_gimple_omp_offloaded (up->stmt)))
+ && maybe_lookup_decl (decl, up))
break;
if (up)
{
tree c;
- for (c = gimple_omp_taskreg_clauses (up->stmt);
- c; c = OMP_CLAUSE_CHAIN (c))
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
- && OMP_CLAUSE_DECL (c) == decl)
- break;
+ if (gimple_code (up->stmt) == GIMPLE_OMP_TARGET)
+ {
+ for (c = gimple_omp_target_clauses (up->stmt);
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_DECL (c) == decl)
+ break;
+ }
+ else
+ for (c = gimple_omp_taskreg_clauses (up->stmt);
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
+ && OMP_CLAUSE_DECL (c) == decl)
+ break;
if (c)
goto maybe_mark_addressable_and_ret;
it's address. But we don't need to take address of privatizations
from that var. */
if (TREE_ADDRESSABLE (var)
- && task_shared_vars
- && bitmap_bit_p (task_shared_vars, DECL_UID (var)))
+ && ((task_shared_vars
+ && bitmap_bit_p (task_shared_vars, DECL_UID (var)))
+ || (global_nonaddressable_vars
+ && bitmap_bit_p (global_nonaddressable_vars, DECL_UID (var)))))
TREE_ADDRESSABLE (copy) = 0;
ctx->block_vars = copy;
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_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD)
+ || ctx->loop_p
|| (code == OMP_CLAUSE_PRIVATE
&& (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
|| gimple_code (ctx->stmt) == GIMPLE_OMP_SECTIONS
}
}
else if (outer)
- {
- if (gimple_code (outer->stmt) == GIMPLE_OMP_GRID_BODY)
- {
- outer = outer->outer;
- gcc_assert (outer
- && gimple_code (outer->stmt) != GIMPLE_OMP_GRID_BODY);
- }
- x = lookup_decl (var, outer);
- }
+ x = lookup_decl (var, outer);
else if (omp_is_reference (var))
/* This can happen with orphaned constructs. If var is reference, it is
possible it is shared and as such valid. */
tree field, type, sfield = NULL_TREE;
splay_tree_key key = (splay_tree_key) var;
+ if ((mask & 16) != 0)
+ {
+ key = (splay_tree_key) &DECL_NAME (var);
+ gcc_checking_assert (key != (splay_tree_key) var);
+ }
if ((mask & 8) != 0)
{
key = (splay_tree_key) &DECL_UID (var);
|| !is_gimple_omp_oacc (ctx->stmt));
type = TREE_TYPE (var);
+ if ((mask & 16) != 0)
+ type = lang_hooks.decls.omp_array_data (var, true);
+
/* Prevent redeclaring the var in the split-off function with a restrict
pointer type. Note that we only clear type itself, restrict qualifiers in
the pointed-to type will be ignored by points-to analysis. */
}
else if (by_ref)
type = build_pointer_type (type);
- else if ((mask & 3) == 1 && omp_is_reference (var))
+ else if ((mask & (32 | 3)) == 1 && omp_is_reference (var))
type = TREE_TYPE (type);
field = build_decl (DECL_SOURCE_LOCATION (var),
side effect of making dwarf2out ignore this member, so for helpful
debugging we clear it later in delete_omp_context. */
DECL_ABSTRACT_ORIGIN (field) = var;
- if (type == TREE_TYPE (var))
+ if ((mask & 16) == 0 && type == TREE_TYPE (var))
{
SET_DECL_ALIGN (field, DECL_ALIGN (var));
DECL_USER_ALIGN (field) = DECL_USER_ALIGN (var);
}
delete ctx->lastprivate_conditional_map;
+ delete ctx->allocate_map;
XDELETE (ctx);
}
tree c, decl;
bool scan_array_reductions = false;
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ALLOCATE
+ && (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) == NULL_TREE
+ /* omp_default_mem_alloc is 1 */
+ || !integer_onep (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))))
+ {
+ if (ctx->allocate_map == NULL)
+ ctx->allocate_map = new hash_map<tree, tree>;
+ ctx->allocate_map->put (OMP_CLAUSE_DECL (c),
+ OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)
+ ? OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)
+ : integer_zero_node);
+ }
+
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
bool by_ref;
case OMP_CLAUSE_SHARED:
decl = OMP_CLAUSE_DECL (c);
+ if (ctx->allocate_map && ctx->allocate_map->get (decl))
+ ctx->allocate_map->remove (decl);
/* Ignore shared directives in teams construct inside of
target construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS
goto do_private;
case OMP_CLAUSE_REDUCTION:
+ /* Collect 'reduction' clauses on OpenACC compute construct. */
+ if (is_gimple_omp_oacc (ctx->stmt)
+ && is_gimple_omp_offloaded (ctx->stmt))
+ {
+ /* No 'reduction' clauses on OpenACC 'kernels'. */
+ gcc_checking_assert (!is_oacc_kernels (ctx));
+ /* Likewise, on OpenACC 'kernels' decomposed parts. */
+ gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx));
+
+ ctx->local_reduction_clauses
+ = tree_cons (NULL, c, ctx->local_reduction_clauses);
+ }
+ /* FALLTHRU */
+
case OMP_CLAUSE_IN_REDUCTION:
decl = OMP_CLAUSE_DECL (c);
+ if (ctx->allocate_map
+ && ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
+ && (OMP_CLAUSE_REDUCTION_INSCAN (c)
+ || OMP_CLAUSE_REDUCTION_TASK (c)))
+ || is_task_ctx (ctx)))
+ {
+ /* For now. */
+ if (ctx->allocate_map->get (decl))
+ ctx->allocate_map->remove (decl);
+ }
if (TREE_CODE (decl) == MEM_REF)
{
tree t = TREE_OPERAND (decl, 0);
if (is_variable_sized (decl))
{
if (is_task_ctx (ctx))
- install_var_field (decl, false, 1, ctx);
+ {
+ if (ctx->allocate_map
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ /* For now. */
+ if (ctx->allocate_map->get (decl))
+ ctx->allocate_map->remove (decl);
+ }
+ install_var_field (decl, false, 1, ctx);
+ }
break;
}
else if (is_taskreg_ctx (ctx))
if (is_task_ctx (ctx)
&& (global || by_ref || omp_is_reference (decl)))
{
- install_var_field (decl, false, 1, ctx);
+ if (ctx->allocate_map
+ && ctx->allocate_map->get (decl))
+ install_var_field (decl, by_ref, 32 | 1, ctx);
+ else
+ install_var_field (decl, false, 1, ctx);
if (!global)
install_var_field (decl, by_ref, 2, ctx);
}
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_USE_DEVICE_ADDR:
decl = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- install_var_field (decl, true, 3, ctx);
+
+ /* Fortran array descriptors. */
+ if (lang_hooks.decls.omp_array_data (decl, true))
+ install_var_field (decl, false, 19, ctx);
+ else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ && !omp_is_reference (decl)
+ && !omp_is_allocatable_or_ptr (decl))
+ || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ install_var_field (decl, true, 11, ctx);
else
- install_var_field (decl, false, 3, ctx);
+ install_var_field (decl, false, 11, ctx);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
&& DECL_P (decl)
&& ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable
&& !lookup_attribute ("omp declare target link",
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
break;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && DECL_P (decl)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (ctx->stmt))
+ {
+ /* If this is an offloaded region, an attach operation should
+ only exist when the pointer variable is mapped in a prior
+ clause. */
+ if (is_gimple_omp_offloaded (ctx->stmt))
+ gcc_assert
+ (maybe_lookup_decl (decl, ctx)
+ || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl))));
+
+ /* By itself, attach/detach is generated as part of pointer
+ variable mapping and should not create new variables in the
+ offloaded region, however sender refs for it must be created
+ for its address to be passed to the runtime. */
+ tree field
+ = build_decl (OMP_CLAUSE_LOCATION (c),
+ FIELD_DECL, NULL_TREE, ptr_type_node);
+ SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+ insert_field_into_struct (ctx->record_type, field);
+ /* To not clash with a map of the pointer variable itself,
+ attach/detach maps have their field looked up by the *clause*
+ tree expression, not the decl. */
+ gcc_assert (!splay_tree_lookup (ctx->field_map,
+ (splay_tree_key) c));
+ splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+ (splay_tree_value) field);
+ break;
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
}
break;
- case OMP_CLAUSE__GRIDDIM_:
- if (ctx->outer)
- {
- scan_omp_op (&OMP_CLAUSE__GRIDDIM__SIZE (c), ctx->outer);
- scan_omp_op (&OMP_CLAUSE__GRIDDIM__GROUP (c), ctx->outer);
- }
+ case OMP_CLAUSE_ORDER:
+ ctx->order_concurrent = true;
+ break;
+
+ case OMP_CLAUSE_BIND:
+ ctx->loop_p = true;
break;
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_TASK_REDUCTION:
+ case OMP_CLAUSE_ALLOCATE:
break;
case OMP_CLAUSE_ALIGNED:
install_var_local (decl, ctx);
}
else if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && (gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD
&& !OMP_CLAUSE__CONDTEMP__ITER (c))
install_var_local (decl, ctx);
break;
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
break;
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (ctx->stmt)
+ && !is_gimple_omp_offloaded (ctx->stmt))
+ break;
if (DECL_P (decl))
{
if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
case OMP_CLAUSE_SIMDLEN:
case OMP_CLAUSE_ALIGNED:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE_ALLOCATE:
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE__REDUCTEMP_:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_NOGROUP:
case OMP_CLAUSE_DEFAULTMAP:
+ case OMP_CLAUSE_ORDER:
+ case OMP_CLAUSE_BIND:
case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_NONTEMPORAL:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
- case OMP_CLAUSE__GRIDDIM_:
case OMP_CLAUSE__SIMT_:
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
GIMPLE_OMP_FOR, add one more temporaries for the total number
of iterations (product of count1 ... countN-1). */
if (omp_find_clause (gimple_omp_for_clauses (for_stmt),
- OMP_CLAUSE_LASTPRIVATE))
- count++;
- else if (msk == GF_OMP_FOR_KIND_FOR
- && omp_find_clause (gimple_omp_parallel_clauses (stmt),
- OMP_CLAUSE_LASTPRIVATE))
- count++;
+ OMP_CLAUSE_LASTPRIVATE)
+ || (msk == GF_OMP_FOR_KIND_FOR
+ && omp_find_clause (gimple_omp_parallel_clauses (stmt),
+ OMP_CLAUSE_LASTPRIVATE)))
+ {
+ 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_taskreg_clauses (stmt);
+ gimple_omp_taskreg_set_clauses (stmt, c);
+ }
+ if (fd.non_rect
+ && fd.last_nonrect == fd.first_nonrect + 1)
+ if (tree v = gimple_omp_for_index (for_stmt, fd.last_nonrect))
+ if (!TYPE_UNSIGNED (TREE_TYPE (v)))
+ {
+ v = gimple_omp_for_index (for_stmt, fd.first_nonrect);
+ tree type2 = TREE_TYPE (v);
+ count++;
+ for (i = 0; i < 3; i++)
+ {
+ tree temp = create_tmp_var (type2);
+ 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_taskreg_clauses (stmt);
+ gimple_omp_taskreg_set_clauses (stmt, c);
+ }
+ }
}
for (i = 0; i < count; i++)
{
DECL_NAMELESS (name) = 1;
TYPE_NAME (ctx->record_type) = name;
TYPE_ARTIFICIAL (ctx->record_type) = 1;
- if (!gimple_omp_parallel_grid_phony (stmt))
- {
- create_omp_child_function (ctx, false);
- gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn);
- }
+ create_omp_child_function (ctx, false);
+ gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn);
scan_sharing_clauses (gimple_omp_parallel_clauses (stmt), ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
return ctx;
}
-/* Return true if ctx is part of an oacc kernels region. */
+/* Return whether CTX's parent compute construct is an OpenACC 'kernels'
+ construct.
+ (This doesn't include OpenACC 'kernels' decomposed parts.) */
static bool
ctx_in_oacc_kernels_region (omp_context *ctx)
return false;
}
-/* Check the parallelism clauses inside a kernels regions.
+/* Check the parallelism clauses inside a OpenACC 'kernels' region.
+ (This doesn't include OpenACC 'kernels' decomposed parts.)
Until kernels handling moves to use the same loop indirection
scheme as parallel, we need to do this checking early. */
{
omp_context *tgt = enclosing_target_ctx (outer_ctx);
- if (!tgt || is_oacc_parallel (tgt))
+ if (!(tgt && is_oacc_kernels (tgt)))
for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
- char const *check = NULL;
-
+ tree c_op0;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_GANG:
- check = "gang";
+ c_op0 = OMP_CLAUSE_GANG_EXPR (c);
break;
case OMP_CLAUSE_WORKER:
- check = "worker";
+ c_op0 = OMP_CLAUSE_WORKER_EXPR (c);
break;
case OMP_CLAUSE_VECTOR:
- check = "vector";
+ c_op0 = OMP_CLAUSE_VECTOR_EXPR (c);
break;
default:
- break;
+ continue;
}
- if (check && OMP_CLAUSE_OPERAND (c, 0))
- error_at (gimple_location (stmt),
- "argument not permitted on %qs clause in"
- " OpenACC %<parallel%>", check);
+ if (c_op0)
+ {
+ /* By construction, this is impossible for OpenACC 'kernels'
+ decomposed parts. */
+ gcc_assert (!(tgt && is_oacc_kernels_decomposed_part (tgt)));
+
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "argument not permitted on %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ if (tgt)
+ inform (gimple_location (tgt->stmt),
+ "enclosing parent compute construct");
+ else if (oacc_get_fn_attrib (current_function_decl))
+ inform (DECL_SOURCE_LOCATION (current_function_decl),
+ "enclosing routine");
+ else
+ gcc_unreachable ();
+ }
}
+ if (tgt && is_oacc_kernels (tgt))
+ check_oacc_kernel_gwv (stmt, ctx);
+
+ /* Collect all variables named in reductions on this loop. Ensure
+ that, if this loop has a reduction on some variable v, and there is
+ a reduction on v somewhere in an outer context, then there is a
+ reduction on v on all intervening loops as well. */
+ tree local_reduction_clauses = NULL;
+ for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ local_reduction_clauses
+ = tree_cons (NULL, c, local_reduction_clauses);
+ }
+ if (ctx->outer_reduction_clauses == NULL && ctx->outer != NULL)
+ ctx->outer_reduction_clauses
+ = chainon (unshare_expr (ctx->outer->local_reduction_clauses),
+ ctx->outer->outer_reduction_clauses);
+ tree outer_reduction_clauses = ctx->outer_reduction_clauses;
+ tree local_iter = local_reduction_clauses;
+ for (; local_iter; local_iter = TREE_CHAIN (local_iter))
+ {
+ tree local_clause = TREE_VALUE (local_iter);
+ tree local_var = OMP_CLAUSE_DECL (local_clause);
+ tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause);
+ bool have_outer_reduction = false;
+ tree ctx_iter = outer_reduction_clauses;
+ for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter))
+ {
+ tree outer_clause = TREE_VALUE (ctx_iter);
+ tree outer_var = OMP_CLAUSE_DECL (outer_clause);
+ tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause);
+ if (outer_var == local_var && outer_op != local_op)
+ {
+ warning_at (OMP_CLAUSE_LOCATION (local_clause), 0,
+ "conflicting reduction operations for %qE",
+ local_var);
+ inform (OMP_CLAUSE_LOCATION (outer_clause),
+ "location of the previous reduction for %qE",
+ outer_var);
+ }
+ if (outer_var == local_var)
+ {
+ have_outer_reduction = true;
+ break;
+ }
+ }
+ if (have_outer_reduction)
+ {
+ /* There is a reduction on outer_var both on this loop and on
+ some enclosing loop. Walk up the context tree until such a
+ loop with a reduction on outer_var is found, and complain
+ about all intervening loops that do not have such a
+ reduction. */
+ struct omp_context *curr_loop = ctx->outer;
+ bool found = false;
+ while (curr_loop != NULL)
+ {
+ tree curr_iter = curr_loop->local_reduction_clauses;
+ for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter))
+ {
+ tree curr_clause = TREE_VALUE (curr_iter);
+ tree curr_var = OMP_CLAUSE_DECL (curr_clause);
+ if (curr_var == local_var)
+ {
+ found = true;
+ break;
+ }
+ }
+ if (!found)
+ warning_at (gimple_location (curr_loop->stmt), 0,
+ "nested loop in reduction needs "
+ "reduction clause for %qE",
+ local_var);
+ else
+ break;
+ curr_loop = curr_loop->outer;
+ }
+ }
+ }
+ ctx->local_reduction_clauses = local_reduction_clauses;
+ ctx->outer_reduction_clauses
+ = chainon (unshare_expr (ctx->local_reduction_clauses),
+ ctx->outer_reduction_clauses);
+
if (tgt && is_oacc_kernels (tgt))
{
/* Strip out reductions, as they are not handled yet. */
while (tree probe = *prev_ptr)
{
tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
-
+
if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
*prev_ptr = *next_ptr;
else
}
gimple_omp_for_set_clauses (stmt, clauses);
- check_oacc_kernel_gwv (stmt, ctx);
}
}
{
tree c;
- if (ctx && gimple_code (ctx->stmt) == GIMPLE_OMP_GRID_BODY)
- /* GRID_BODY is an artificial construct, nesting rules will be checked in
- the original copy of its contents. */
- return true;
-
/* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
inside an OpenACC CTX. */
if (!(is_gimple_omp (stmt)
&& gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
ctx = ctx->outer;
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD
+ && !ctx->loop_p)
{
c = NULL_TREE;
+ if (ctx->order_concurrent
+ && (gimple_code (stmt) == GIMPLE_OMP_ORDERED
+ || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD
+ || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE))
+ {
+ error_at (gimple_location (stmt),
+ "OpenMP constructs other than %<parallel%>, %<loop%>"
+ " or %<simd%> may not be nested inside a region with"
+ " the %<order(concurrent)%> clause");
+ return false;
+ }
if (gimple_code (stmt) == GIMPLE_OMP_ORDERED)
{
c = gimple_omp_ordered_clauses (as_a <gomp_ordered *> (stmt));
{
error_at (gimple_location (stmt),
"%<ordered simd threads%> must be closely "
- "nested inside of %<for simd%> region");
+ "nested inside of %<%s simd%> region",
+ lang_GNU_Fortran () ? "do" : "for");
return false;
}
return true;
|| gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE
|| gimple_code (stmt) == GIMPLE_OMP_SCAN)
return true;
+ else if (gimple_code (stmt) == GIMPLE_OMP_FOR
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD)
+ return true;
error_at (gimple_location (stmt),
- "OpenMP constructs other than %<#pragma omp ordered simd%>"
- " or %<#pragma omp atomic%> may not be nested inside"
- " %<simd%> region");
+ "OpenMP constructs other than "
+ "%<ordered simd%>, %<simd%>, %<loop%> or %<atomic%> may "
+ "not be nested inside %<simd%> region");
return false;
}
else if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
{
if ((gimple_code (stmt) != GIMPLE_OMP_FOR
- || ((gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_DISTRIBUTE)
- && (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_GRID_LOOP)))
+ || (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_DISTRIBUTE
+ && omp_find_clause (gimple_omp_for_clauses (stmt),
+ OMP_CLAUSE_BIND) == NULL_TREE))
&& gimple_code (stmt) != GIMPLE_OMP_PARALLEL)
{
error_at (gimple_location (stmt),
- "only %<distribute%> or %<parallel%> regions are "
- "allowed to be strictly nested inside %<teams%> "
- "region");
+ "only %<distribute%>, %<parallel%> or %<loop%> "
+ "regions are allowed to be strictly nested inside "
+ "%<teams%> region");
return false;
}
}
+ else if (ctx->order_concurrent
+ && gimple_code (stmt) != GIMPLE_OMP_PARALLEL
+ && (gimple_code (stmt) != GIMPLE_OMP_FOR
+ || gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_SIMD)
+ && gimple_code (stmt) != GIMPLE_OMP_SCAN)
+ {
+ if (ctx->loop_p)
+ error_at (gimple_location (stmt),
+ "OpenMP constructs other than %<parallel%>, %<loop%> or "
+ "%<simd%> may not be nested inside a %<loop%> region");
+ else
+ error_at (gimple_location (stmt),
+ "OpenMP constructs other than %<parallel%>, %<loop%> or "
+ "%<simd%> may not be nested inside a region with "
+ "the %<order(concurrent)%> clause");
+ return false;
+ }
}
switch (gimple_code (stmt))
{
case GIMPLE_OMP_FOR:
- if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD)
+ if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_SIMD)
return true;
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
{
/* We split taskloop into task and nested taskloop in it. */
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP)
return true;
+ /* For now, hope this will change and loop bind(parallel) will not
+ be allowed in lots of contexts. */
+ if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR
+ && omp_find_clause (gimple_omp_for_clauses (stmt), OMP_CLAUSE_BIND))
+ return true;
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
{
bool ok = false;
{
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
ok = true;
break;
const char *construct
= (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt))
== BUILT_IN_GOMP_CANCEL)
- ? "#pragma omp cancel"
- : "#pragma omp cancellation point";
+ ? "cancel"
+ : "cancellation point";
if (ctx == NULL)
{
error_at (gimple_location (stmt), "orphaned %qs construct",
{
case 1:
if (gimple_code (ctx->stmt) != GIMPLE_OMP_PARALLEL)
- bad = "#pragma omp parallel";
+ bad = "parallel";
else if (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt))
== BUILT_IN_GOMP_CANCEL
&& !integer_zerop (gimple_call_arg (stmt, 1)))
case 2:
if (gimple_code (ctx->stmt) != GIMPLE_OMP_FOR
|| gimple_omp_for_kind (ctx->stmt) != GF_OMP_FOR_KIND_FOR)
- bad = "#pragma omp for";
+ bad = "for";
else if (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt))
== BUILT_IN_GOMP_CANCEL
&& !integer_zerop (gimple_call_arg (stmt, 1)))
if (omp_find_clause (gimple_omp_for_clauses (ctx->stmt),
OMP_CLAUSE_NOWAIT))
warning_at (gimple_location (stmt), 0,
- "%<#pragma omp cancel for%> inside "
+ "%<cancel for%> inside "
"%<nowait%> for construct");
if (omp_find_clause (gimple_omp_for_clauses (ctx->stmt),
OMP_CLAUSE_ORDERED))
warning_at (gimple_location (stmt), 0,
- "%<#pragma omp cancel for%> inside "
+ "%<cancel for%> inside "
"%<ordered%> for construct");
}
kind = "for";
case 4:
if (gimple_code (ctx->stmt) != GIMPLE_OMP_SECTIONS
&& gimple_code (ctx->stmt) != GIMPLE_OMP_SECTION)
- bad = "#pragma omp sections";
+ bad = "sections";
else if (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt))
== BUILT_IN_GOMP_CANCEL
&& !integer_zerop (gimple_call_arg (stmt, 1)))
(ctx->stmt),
OMP_CLAUSE_NOWAIT))
warning_at (gimple_location (stmt), 0,
- "%<#pragma omp cancel sections%> inside "
+ "%<cancel sections%> inside "
"%<nowait%> sections construct");
}
else
(ctx->outer->stmt),
OMP_CLAUSE_NOWAIT))
warning_at (gimple_location (stmt), 0,
- "%<#pragma omp cancel sections%> inside "
+ "%<cancel sections%> inside "
"%<nowait%> sections construct");
}
}
&& (!is_taskloop_ctx (ctx)
|| ctx->outer == NULL
|| !is_task_ctx (ctx->outer)))
- bad = "#pragma omp task";
+ bad = "task";
else
{
for (omp_context *octx = ctx->outer;
return true;
error_at (gimple_location (stmt),
"barrier region may not be closely nested inside "
- "of work-sharing, %<critical%>, %<ordered%>, "
- "%<master%>, explicit %<task%> or %<taskloop%> "
- "region");
+ "of work-sharing, %<loop%>, %<critical%>, "
+ "%<ordered%>, %<master%>, explicit %<task%> or "
+ "%<taskloop%> region");
return false;
}
error_at (gimple_location (stmt),
"work-sharing region may not be closely nested inside "
- "of work-sharing, %<critical%>, %<ordered%>, "
+ "of work-sharing, %<loop%>, %<critical%>, %<ordered%>, "
"%<master%>, explicit %<task%> or %<taskloop%> region");
return false;
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
error_at (gimple_location (stmt),
"%<master%> region may not be closely nested inside "
- "of work-sharing, explicit %<task%> or %<taskloop%> "
- "region");
+ "of work-sharing, %<loop%>, explicit %<task%> or "
+ "%<taskloop%> region");
return false;
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TEAMS:
stmt_name = "target exit data"; 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_SERIAL: stmt_name = "serial"; 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:
case GF_OMP_TARGET_KIND_OACC_DECLARE: stmt_name = "declare"; break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
+ case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
+ /* OpenACC 'kernels' decomposed parts. */
+ stmt_name = "kernels"; break;
default: gcc_unreachable ();
}
switch (gimple_omp_target_kind (ctx->stmt))
ctx_stmt_name = "parallel"; break;
case GF_OMP_TARGET_KIND_OACC_KERNELS:
ctx_stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
+ ctx_stmt_name = "serial"; break;
case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
ctx_stmt_name = "host_data"; break;
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
+ case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
+ /* OpenACC 'kernels' decomposed parts. */
+ ctx_stmt_name = "kernels"; break;
default: gcc_unreachable ();
}
return true;
tree declname = DECL_NAME (fndecl);
- if (!declname)
+ if (!declname
+ || (DECL_CONTEXT (fndecl) != NULL_TREE
+ && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL)
+ || !TREE_PUBLIC (fndecl))
return false;
+
const char *name = IDENTIFIER_POINTER (declname);
return !strcmp (name, "setjmp") || !strcmp (name, "longjmp");
}
+/* Return true if FNDECL is an omp_* runtime API call. */
+
+static bool
+omp_runtime_api_call (const_tree fndecl)
+{
+ tree declname = DECL_NAME (fndecl);
+ if (!declname
+ || (DECL_CONTEXT (fndecl) != NULL_TREE
+ && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL)
+ || !TREE_PUBLIC (fndecl))
+ return false;
+
+ const char *name = IDENTIFIER_POINTER (declname);
+ if (strncmp (name, "omp_", 4) != 0)
+ return false;
+
+ static const char *omp_runtime_apis[] =
+ {
+ /* This array has 3 sections. First omp_* calls that don't
+ have any suffixes. */
+ "target_alloc",
+ "target_associate_ptr",
+ "target_disassociate_ptr",
+ "target_free",
+ "target_is_present",
+ "target_memcpy",
+ "target_memcpy_rect",
+ NULL,
+ /* Now omp_* calls that are available as omp_* and omp_*_. */
+ "capture_affinity",
+ "destroy_lock",
+ "destroy_nest_lock",
+ "display_affinity",
+ "get_active_level",
+ "get_affinity_format",
+ "get_cancellation",
+ "get_default_device",
+ "get_dynamic",
+ "get_initial_device",
+ "get_level",
+ "get_max_active_levels",
+ "get_max_task_priority",
+ "get_max_threads",
+ "get_nested",
+ "get_num_devices",
+ "get_num_places",
+ "get_num_procs",
+ "get_num_teams",
+ "get_num_threads",
+ "get_partition_num_places",
+ "get_place_num",
+ "get_proc_bind",
+ "get_team_num",
+ "get_thread_limit",
+ "get_thread_num",
+ "get_wtick",
+ "get_wtime",
+ "in_final",
+ "in_parallel",
+ "init_lock",
+ "init_nest_lock",
+ "is_initial_device",
+ "pause_resource",
+ "pause_resource_all",
+ "set_affinity_format",
+ "set_lock",
+ "set_nest_lock",
+ "test_lock",
+ "test_nest_lock",
+ "unset_lock",
+ "unset_nest_lock",
+ NULL,
+ /* And finally calls available as omp_*, omp_*_ and omp_*_8_. */
+ "get_ancestor_thread_num",
+ "get_partition_place_nums",
+ "get_place_num_procs",
+ "get_place_proc_ids",
+ "get_schedule",
+ "get_team_size",
+ "set_default_device",
+ "set_dynamic",
+ "set_max_active_levels",
+ "set_nested",
+ "set_num_threads",
+ "set_schedule"
+ };
+
+ int mode = 0;
+ for (unsigned i = 0; i < ARRAY_SIZE (omp_runtime_apis); i++)
+ {
+ if (omp_runtime_apis[i] == NULL)
+ {
+ mode++;
+ continue;
+ }
+ size_t len = strlen (omp_runtime_apis[i]);
+ if (strncmp (name + 4, omp_runtime_apis[i], len) == 0
+ && (name[4 + len] == '\0'
+ || (mode > 0
+ && name[4 + len] == '_'
+ && (name[4 + len + 1] == '\0'
+ || (mode > 1
+ && strcmp (name + 4 + len + 1, "8_") == 0)))))
+ return true;
+ }
+ return false;
+}
/* Helper function for scan_omp.
tree fndecl = gimple_call_fndecl (stmt);
if (fndecl)
{
- if (setjmp_or_longjmp_p (fndecl)
- && ctx
+ if (ctx
&& gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD
+ && setjmp_or_longjmp_p (fndecl)
+ && !ctx->loop_p)
{
remove = true;
error_at (gimple_location (stmt),
- "setjmp/longjmp inside simd construct");
+ "setjmp/longjmp inside %<simd%> construct");
}
else if (DECL_BUILT_IN_CLASS (fndecl) == BUILT_IN_NORMAL)
switch (DECL_FUNCTION_CODE (fndecl))
default:
break;
}
+ else if (ctx)
+ {
+ omp_context *octx = ctx;
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_SCAN && ctx->outer)
+ octx = ctx->outer;
+ if (octx->order_concurrent && omp_runtime_api_call (fndecl))
+ {
+ remove = true;
+ error_at (gimple_location (stmt),
+ "OpenMP runtime API call %qD in a region with "
+ "%<order(concurrent)%> clause", fndecl);
+ }
+ }
}
}
if (remove)
if ((gimple_omp_for_kind (as_a <gomp_for *> (stmt))
== GF_OMP_FOR_KIND_SIMD)
&& omp_maybe_offloaded_ctx (ctx)
- && omp_max_simt_vf ())
+ && omp_max_simt_vf ()
+ && gimple_omp_for_collapse (stmt) == 1)
scan_omp_simd (gsi, as_a <gomp_for *> (stmt), ctx);
else
scan_omp_for (as_a <gomp_for *> (stmt), ctx);
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
- case GIMPLE_OMP_GRID_BODY:
ctx = new_omp_context (stmt, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
break;
break;
case GIMPLE_OMP_TARGET:
- scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+ if (is_gimple_omp_offloaded (stmt))
+ {
+ taskreg_nesting_level++;
+ scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+ taskreg_nesting_level--;
+ }
+ else
+ scan_omp_target (as_a <gomp_target *> (stmt), ctx);
break;
case GIMPLE_OMP_TEAMS:
/* Otherwise return implementation defined alignment. */
unsigned int al = 1;
opt_scalar_mode mode_iter;
- auto_vector_sizes sizes;
- targetm.vectorize.autovectorize_vector_sizes (&sizes, true);
- poly_uint64 vs = 0;
- for (unsigned int i = 0; i < sizes.length (); ++i)
- vs = ordered_max (vs, sizes[i]);
+ auto_vector_modes modes;
+ targetm.vectorize.autovectorize_vector_modes (&modes, true);
static enum mode_class classes[]
= { MODE_INT, MODE_VECTOR_INT, MODE_FLOAT, MODE_VECTOR_FLOAT };
for (int i = 0; i < 4; i += 2)
machine_mode vmode = targetm.vectorize.preferred_simd_mode (mode);
if (GET_MODE_CLASS (vmode) != classes[i + 1])
continue;
- while (maybe_ne (vs, 0U)
- && known_lt (GET_MODE_SIZE (vmode), vs)
- && GET_MODE_2XWIDER_MODE (vmode).exists ())
- vmode = GET_MODE_2XWIDER_MODE (vmode).require ();
+ machine_mode alt_vmode;
+ for (unsigned int j = 0; j < modes.length (); ++j)
+ if (related_vector_mode (modes[j], mode).exists (&alt_vmode)
+ && known_ge (GET_MODE_SIZE (alt_vmode), GET_MODE_SIZE (vmode)))
+ vmode = alt_vmode;
tree type = lang_hooks.types.type_for_mode (mode, 1);
if (type == NULL_TREE || TYPE_MODE (type) != mode)
continue;
- poly_uint64 nelts = exact_div (GET_MODE_SIZE (vmode),
- GET_MODE_SIZE (mode));
- type = build_vector_type (type, nelts);
+ type = build_vector_type_for_mode (type, vmode);
if (TYPE_MODE (type) != vmode)
continue;
if (TYPE_ALIGN_UNIT (type) > al)
/* This structure is part of the interface between lower_rec_simd_input_clauses
and lower_rec_input_clauses. */
-struct omplow_simd_context {
+class omplow_simd_context {
+public:
omplow_simd_context () { memset (this, 0, sizeof (*this)); }
tree idx;
tree lane;
DECL_ATTRIBUTES (ivar) = tree_cons (get_identifier ("omp simt private"),
NULL, DECL_ATTRIBUTES (ivar));
sctx->simt_eargs.safe_push (build1 (ADDR_EXPR, ptype, ivar));
- tree clobber = build_constructor (type, NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (type);
gimple *g = gimple_build_assign (ivar, clobber);
gimple_seq_add_stmt (&sctx->simt_dlist, g);
}
return v;
}
+/* Lower early initialization of privatized variable NEW_VAR
+ if it needs an allocator (has allocate clause). */
+
+static bool
+lower_private_allocate (tree var, tree new_var, tree &allocator,
+ tree &allocate_ptr, gimple_seq *ilist,
+ omp_context *ctx, bool is_ref, tree size)
+{
+ if (allocator)
+ return false;
+ gcc_assert (allocate_ptr == NULL_TREE);
+ if (ctx->allocate_map
+ && (DECL_P (new_var) || (TYPE_P (new_var) && size)))
+ if (tree *allocatorp = ctx->allocate_map->get (var))
+ allocator = *allocatorp;
+ if (allocator == NULL_TREE)
+ return false;
+ if (!is_ref && omp_is_reference (var))
+ {
+ allocator = NULL_TREE;
+ return false;
+ }
+
+ if (TREE_CODE (allocator) != INTEGER_CST)
+ allocator = build_outer_var_ref (allocator, ctx);
+ allocator = fold_convert (pointer_sized_int_node, allocator);
+ if (TREE_CODE (allocator) != INTEGER_CST)
+ {
+ tree var = create_tmp_var (TREE_TYPE (allocator));
+ gimplify_assign (var, allocator, ilist);
+ allocator = var;
+ }
+
+ tree ptr_type, align, sz = size;
+ if (TYPE_P (new_var))
+ {
+ ptr_type = build_pointer_type (new_var);
+ align = build_int_cst (size_type_node, TYPE_ALIGN_UNIT (new_var));
+ }
+ else if (is_ref)
+ {
+ ptr_type = build_pointer_type (TREE_TYPE (TREE_TYPE (new_var)));
+ align = build_int_cst (size_type_node,
+ TYPE_ALIGN_UNIT (TREE_TYPE (ptr_type)));
+ }
+ else
+ {
+ ptr_type = build_pointer_type (TREE_TYPE (new_var));
+ align = build_int_cst (size_type_node, DECL_ALIGN_UNIT (new_var));
+ if (sz == NULL_TREE)
+ sz = fold_convert (size_type_node, DECL_SIZE_UNIT (new_var));
+ }
+ if (TREE_CODE (sz) != INTEGER_CST)
+ {
+ tree szvar = create_tmp_var (size_type_node);
+ gimplify_assign (szvar, sz, ilist);
+ sz = szvar;
+ }
+ allocate_ptr = create_tmp_var (ptr_type);
+ tree a = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC);
+ gimple *g = gimple_build_call (a, 3, align, sz, allocator);
+ gimple_call_set_lhs (g, allocate_ptr);
+ gimple_seq_add_stmt (ilist, g);
+ if (!is_ref)
+ {
+ tree x = build_simple_mem_ref (allocate_ptr);
+ TREE_THIS_NOTRAP (x) = 1;
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ return true;
+}
+
/* 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_SIMD);
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD);
omplow_simd_context sctx = omplow_simd_context ();
tree simt_lane = NULL_TREE, simtrec = NULL_TREE;
tree ivar = NULL_TREE, lvar = NULL_TREE, uid = NULL_TREE;
bool task_reduction_p = false;
bool task_reduction_needs_orig_p = false;
tree cond = NULL_TREE;
+ tree allocator, allocate_ptr;
switch (c_kind)
{
if (task_reduction_p != (pass >= 2))
continue;
+ allocator = NULL_TREE;
+ allocate_ptr = NULL_TREE;
new_var = var = OMP_CLAUSE_DECL (c);
if ((c_kind == OMP_CLAUSE_REDUCTION
|| c_kind == OMP_CLAUSE_IN_REDUCTION)
tree type = TREE_TYPE (d);
gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
+ tree sz = v;
const char *name = get_name (orig_var);
+ if (pass != 3 && !TREE_CONSTANT (v))
+ {
+ tree t = maybe_lookup_decl (v, ctx);
+ if (t)
+ v = t;
+ else
+ v = maybe_lookup_decl_in_outer_ctx (v, ctx);
+ gimplify_expr (&v, ilist, NULL, is_gimple_val, fb_rvalue);
+ t = fold_build2_loc (clause_loc, PLUS_EXPR,
+ TREE_TYPE (v), v,
+ build_int_cst (TREE_TYPE (v), 1));
+ sz = fold_build2_loc (clause_loc, MULT_EXPR,
+ TREE_TYPE (v), t,
+ TYPE_SIZE_UNIT (TREE_TYPE (type)));
+ }
if (pass == 3)
{
tree xv = create_tmp_var (ptr_type_node);
gimplify_assign (cond, x, ilist);
x = xv;
}
+ else if (lower_private_allocate (var, type, allocator,
+ allocate_ptr, ilist, ctx,
+ true,
+ TREE_CONSTANT (v)
+ ? TYPE_SIZE_UNIT (type)
+ : sz))
+ x = allocate_ptr;
else if (TREE_CONSTANT (v))
{
x = create_tmp_var_raw (type, name);
{
tree atmp
= builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
- tree t = maybe_lookup_decl (v, ctx);
- if (t)
- v = t;
- else
- v = maybe_lookup_decl_in_outer_ctx (v, ctx);
- gimplify_expr (&v, ilist, NULL, is_gimple_val, fb_rvalue);
- t = fold_build2_loc (clause_loc, PLUS_EXPR,
- TREE_TYPE (v), v,
- build_int_cst (TREE_TYPE (v), 1));
- t = fold_build2_loc (clause_loc, MULT_EXPR,
- TREE_TYPE (v), t,
- TYPE_SIZE_UNIT (TREE_TYPE (type)));
tree al = size_int (TYPE_ALIGN (TREE_TYPE (type)));
- x = build_call_expr_loc (clause_loc, atmp, 2, t, al);
+ x = build_call_expr_loc (clause_loc, atmp, 2, sz, al);
}
tree ptype = build_pointer_type (TREE_TYPE (type));
gimple_seq_add_stmt (dlist, g);
gimple_seq_add_stmt (dlist, gimple_build_label (end2));
}
+ if (allocator)
+ {
+ tree f = builtin_decl_explicit (BUILT_IN_GOMP_FREE);
+ g = gimple_build_call (f, 2, allocate_ptr, allocator);
+ gimple_seq_add_stmt (dlist, g);
+ }
continue;
}
else if (pass == 2)
if (c_kind != OMP_CLAUSE_FIRSTPRIVATE || !is_task_ctx (ctx))
{
- gcall *stmt;
- tree tmp, atmp;
+ tree tmp;
ptr = DECL_VALUE_EXPR (new_var);
gcc_assert (TREE_CODE (ptr) == INDIRECT_REF);
gcc_assert (DECL_P (ptr));
x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
- /* void *tmp = __builtin_alloca */
- atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
- stmt = gimple_build_call (atmp, 2, x,
- size_int (DECL_ALIGN (var)));
- tmp = create_tmp_var_raw (ptr_type_node);
- gimple_add_tmp_var (tmp);
- gimple_call_set_lhs (stmt, tmp);
-
- gimple_seq_add_stmt (ilist, stmt);
+ if (lower_private_allocate (var, new_var, allocator,
+ allocate_ptr, ilist, ctx,
+ false, x))
+ tmp = allocate_ptr;
+ else
+ {
+ /* void *tmp = __builtin_alloca */
+ tree atmp
+ = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+ gcall *stmt
+ = gimple_build_call (atmp, 2, x,
+ size_int (DECL_ALIGN (var)));
+ cfun->calls_alloca = 1;
+ tmp = create_tmp_var_raw (ptr_type_node);
+ gimple_add_tmp_var (tmp);
+ gimple_call_set_lhs (stmt, tmp);
+
+ gimple_seq_add_stmt (ilist, stmt);
+ }
x = fold_convert_loc (clause_loc, TREE_TYPE (ptr), tmp);
gimplify_assign (ptr, x, ilist);
if (c_kind == OMP_CLAUSE_FIRSTPRIVATE && is_task_ctx (ctx))
{
x = build_receiver_ref (var, false, ctx);
- x = build_fold_addr_expr_loc (clause_loc, x);
+ if (ctx->allocate_map)
+ if (tree *allocatep = ctx->allocate_map->get (var))
+ {
+ allocator = *allocatep;
+ if (TREE_CODE (allocator) != INTEGER_CST)
+ allocator = build_outer_var_ref (allocator, ctx);
+ allocator = fold_convert (pointer_sized_int_node,
+ allocator);
+ allocate_ptr = unshare_expr (x);
+ }
+ if (allocator == NULL_TREE)
+ x = build_fold_addr_expr_loc (clause_loc, x);
}
+ else if (lower_private_allocate (var, new_var, allocator,
+ allocate_ptr,
+ ilist, ctx, true, x))
+ x = allocate_ptr;
else if (TREE_CONSTANT (x))
{
/* For reduction in SIMD loop, defer adding the
x = NULL;
do_private:
tree nx;
- nx = lang_hooks.decls.omp_clause_default_ctor
- (c, unshare_expr (new_var), x);
+ bool copy_ctor;
+ copy_ctor = false;
+ lower_private_allocate (var, new_var, allocator, allocate_ptr,
+ ilist, ctx, false, NULL_TREE);
+ nx = unshare_expr (new_var);
+ if (is_simd
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
+ && OMP_CLAUSE_LASTPRIVATE_LOOP_IV (c))
+ copy_ctor = true;
+ if (copy_ctor)
+ nx = lang_hooks.decls.omp_clause_copy_ctor (c, nx, x);
+ else
+ nx = lang_hooks.decls.omp_clause_default_ctor (c, nx, x);
if (is_simd)
{
tree y = lang_hooks.decls.omp_clause_dtor (c, new_var);
if ((TREE_ADDRESSABLE (new_var) || nx || y
- || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
+ || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
+ && (gimple_omp_for_collapse (ctx->stmt) != 1
+ || (gimple_omp_for_index (ctx->stmt, 0)
+ != new_var)))
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE__CONDTEMP_
|| omp_is_reference (var))
&& lower_rec_simd_input_clauses (new_var, ctx, &sctx,
}
if (nx)
- x = lang_hooks.decls.omp_clause_default_ctor
- (c, unshare_expr (ivar), x);
+ {
+ tree iv = unshare_expr (ivar);
+ if (copy_ctor)
+ x = lang_hooks.decls.omp_clause_copy_ctor (c, iv,
+ x);
+ else
+ x = lang_hooks.decls.omp_clause_default_ctor (c,
+ iv,
+ x);
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE__CONDTEMP_)
{
x = build2 (MODIFY_EXPR, TREE_TYPE (ivar),
lower_omp (&tseq, ctx->outer);
gimple_seq_add_seq (&llist[1], tseq);
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
+ && ctx->for_simd_scan_phase)
+ {
+ x = unshare_expr (ivar);
+ tree orig_v
+ = build_outer_var_ref (var, ctx,
+ OMP_CLAUSE_LASTPRIVATE);
+ x = lang_hooks.decls.omp_clause_assign_op (c, x,
+ orig_v);
+ gimplify_and_add (x, &llist[0]);
+ }
if (y)
{
y = lang_hooks.decls.omp_clause_dtor (c, ivar);
}
if (nx)
gimplify_and_add (nx, ilist);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
+ && is_simd
+ && ctx->for_simd_scan_phase)
+ {
+ tree orig_v = build_outer_var_ref (var, ctx,
+ OMP_CLAUSE_LASTPRIVATE);
+ x = lang_hooks.decls.omp_clause_assign_op (c, new_var,
+ orig_v);
+ gimplify_and_add (x, ilist);
+ }
/* FALLTHRU */
do_dtor:
x = lang_hooks.decls.omp_clause_dtor (c, new_var);
if (x)
gimplify_and_add (x, dlist);
+ if (allocator)
+ {
+ if (!is_gimple_val (allocator))
+ {
+ tree avar = create_tmp_var (TREE_TYPE (allocator));
+ gimplify_assign (avar, allocator, dlist);
+ allocator = avar;
+ }
+ if (!is_gimple_val (allocate_ptr))
+ {
+ tree apvar = create_tmp_var (TREE_TYPE (allocate_ptr));
+ gimplify_assign (apvar, allocate_ptr, dlist);
+ allocate_ptr = apvar;
+ }
+ tree f = builtin_decl_explicit (BUILT_IN_GOMP_FREE);
+ gimple *g
+ = gimple_build_call (f, 2, allocate_ptr, allocator);
+ gimple_seq_add_stmt (dlist, g);
+ }
break;
case OMP_CLAUSE_LINEAR:
|| use_pointer_for_field (var, NULL))
{
x = build_receiver_ref (var, false, ctx);
+ if (ctx->allocate_map)
+ if (tree *allocatep = ctx->allocate_map->get (var))
+ {
+ allocator = *allocatep;
+ if (TREE_CODE (allocator) != INTEGER_CST)
+ allocator = build_outer_var_ref (allocator, ctx);
+ allocator = fold_convert (pointer_sized_int_node,
+ allocator);
+ allocate_ptr = unshare_expr (x);
+ x = build_simple_mem_ref (x);
+ TREE_THIS_NOTRAP (x) = 1;
+ }
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
goto do_dtor;
goto do_dtor;
}
do_firstprivate:
+ lower_private_allocate (var, new_var, allocator, allocate_ptr,
+ ilist, ctx, false, NULL_TREE);
x = build_outer_var_ref (var, ctx);
if (is_simd)
{
}
else
{
+ lower_private_allocate (var, new_var, allocator,
+ allocate_ptr, ilist, ctx, false,
+ NULL_TREE);
x = build_outer_var_ref (var, ctx);
if (omp_is_reference (var)
}
else
{
+ lower_private_allocate (var, new_var, allocator,
+ allocate_ptr, ilist, ctx,
+ false, NULL_TREE);
if (omp_is_reference (var) && is_simd)
handle_simd_reference (clause_loc, new_vard, ilist);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
ref = build_outer_var_ref (var, ctx);
gimplify_assign (ref, x, dlist);
}
+ if (allocator)
+ goto do_dtor;
}
}
break;
}
if (tskred_avar)
{
- tree clobber = build_constructor (TREE_TYPE (tskred_avar), NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (TREE_TYPE (tskred_avar));
gimple_seq_add_stmt (ilist, gimple_build_assign (tskred_avar, clobber));
}
&& OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (c))
{
tree o = lookup_decl (OMP_CLAUSE_DECL (c), ctx);
- tree *v
- = ctx->lastprivate_conditional_map->get (o);
- tree po = lookup_decl (OMP_CLAUSE_DECL (c), ctx->outer);
- tree *pv
- = ctx->outer->lastprivate_conditional_map->get (po);
+ omp_context *outer = ctx->outer;
+ if (gimple_code (outer->stmt) == GIMPLE_OMP_SCAN)
+ outer = outer->outer;
+ tree *v = ctx->lastprivate_conditional_map->get (o);
+ tree po = lookup_decl (OMP_CLAUSE_DECL (c), outer);
+ tree *pv = outer->lastprivate_conditional_map->get (po);
*v = *pv;
}
}
tree cond_ptr = NULL_TREE;
tree iter_var = NULL_TREE;
bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD);
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD);
tree next = *clauses;
for (tree c = *clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
bool maybe_simt = false;
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD)
{
maybe_simt = omp_find_clause (orig_clauses, OMP_CLAUSE__SIMT_);
simduid = omp_find_clause (orig_clauses, OMP_CLAUSE__SIMDUID_);
x = NULL_TREE;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
- && OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV (c))
+ && OMP_CLAUSE_LASTPRIVATE_LOOP_IV (c)
+ && is_taskloop_ctx (ctx))
{
- gcc_checking_assert (is_taskloop_ctx (ctx));
tree ovar = maybe_lookup_decl_in_outer_ctx (var,
ctx->outer->outer);
if (is_global_var (ovar))
for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
{
+ /* No 'reduction' clauses on OpenACC 'kernels'. */
+ gcc_checking_assert (!is_oacc_kernels (ctx));
+ /* Likewise, on OpenACC 'kernels' decomposed parts. */
+ gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx));
+
tree orig = OMP_CLAUSE_DECL (c);
tree var = maybe_lookup_decl (orig, ctx);
tree ref_to_res = NULL_TREE;
break;
case GIMPLE_OMP_TARGET:
- if (gimple_omp_target_kind (probe->stmt)
- != GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ /* No 'reduction' clauses inside OpenACC 'kernels'
+ regions. */
+ gcc_checking_assert (!is_oacc_kernels (probe));
+
+ if (!is_gimple_omp_offloaded (probe->stmt))
goto do_lookup;
cls = gimple_omp_target_clauses (probe->stmt);
/* 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_SIMD)
+ && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD)
return;
/* inscan reductions are handled elsewhere. */
continue;
nvar = maybe_lookup_decl (ovar, ctx);
- if (!nvar || !DECL_HAS_VALUE_EXPR_P (nvar))
+ if (!nvar
+ || !DECL_HAS_VALUE_EXPR_P (nvar)
+ || (ctx->allocate_map
+ && ctx->allocate_map->get (ovar)))
continue;
/* If CTX is a nested parallel directive. Find the immediately
tag |= OLF_GANG_STATIC;
}
- /* In a parallel region, loops are implicitly INDEPENDENT. */
omp_context *tgt = enclosing_target_ctx (ctx);
- if (!tgt || is_oacc_parallel (tgt))
+ if (!tgt || is_oacc_parallel_or_serial (tgt))
+ ;
+ else if (is_oacc_kernels (tgt))
+ /* Not using this loops handling inside OpenACC 'kernels' regions. */
+ gcc_unreachable ();
+ else if (is_oacc_kernels_decomposed_part (tgt))
+ ;
+ else
+ gcc_unreachable ();
+
+ /* In a parallel region, loops are implicitly INDEPENDENT. */
+ if (!tgt || is_oacc_parallel_or_serial (tgt))
tag |= OLF_INDEPENDENT;
+ /* Loops inside OpenACC 'kernels' decomposed parts' regions are expected to
+ have an explicit 'seq' or 'independent' clause, and no 'auto' clause. */
+ if (tgt && is_oacc_kernels_decomposed_part (tgt))
+ {
+ gcc_assert (tag & (OLF_SEQ | OLF_INDEPENDENT));
+ gcc_assert (!(tag & OLF_AUTO));
+ }
+
if (tag & OLF_TILE)
/* Tiling could use all 3 levels. */
levels = 3;
if (ctx->record_type)
{
gimple_stmt_iterator gsi = gsi_start (bind_body_tail);
- tree clobber = build_constructor (ctx->record_type, NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (ctx->record_type);
gsi_insert_after (&gsi, gimple_build_assign (ctx->sender_decl,
clobber), GSI_SAME_STMT);
}
bool input_phase = has_clauses ^ octx->scan_inclusive;
bool is_simd = (gimple_code (octx->stmt) == GIMPLE_OMP_FOR
- && (gimple_omp_for_kind (octx->stmt) & GF_OMP_FOR_SIMD));
+ && gimple_omp_for_kind (octx->stmt) == GF_OMP_FOR_KIND_SIMD);
bool is_for = (gimple_code (octx->stmt) == GIMPLE_OMP_FOR
&& gimple_omp_for_kind (octx->stmt) == GF_OMP_FOR_KIND_FOR
&& !gimple_omp_for_combined_p (octx->stmt));
cond_code = EQ_EXPR;
}
- if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_GRID_LOOP
- || gimple_omp_for_grid_phony (fd->for_stmt))
- cond = omp_grid_lastprivate_predicate (fd);
- else
+ tree n2 = fd->loop.n2;
+ if (fd->collapse > 1
+ && TREE_CODE (n2) != INTEGER_CST
+ && gimple_omp_for_combined_into_p (fd->for_stmt))
{
- tree n2 = fd->loop.n2;
- if (fd->collapse > 1
- && TREE_CODE (n2) != INTEGER_CST
- && gimple_omp_for_combined_into_p (fd->for_stmt))
+ struct omp_context *taskreg_ctx = NULL;
+ if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
{
- struct omp_context *taskreg_ctx = NULL;
- if (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
+ || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE)
{
- gomp_for *gfor = as_a <gomp_for *> (ctx->outer->stmt);
- if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR
- || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE)
+ if (gimple_omp_for_combined_into_p (gfor))
{
- if (gimple_omp_for_combined_into_p (gfor))
- {
- gcc_assert (ctx->outer->outer
- && is_parallel_ctx (ctx->outer->outer));
- taskreg_ctx = ctx->outer->outer;
- }
- else
- {
- struct omp_for_data outer_fd;
- omp_extract_for_data (gfor, &outer_fd, NULL);
- n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
- }
+ gcc_assert (ctx->outer->outer
+ && is_parallel_ctx (ctx->outer->outer));
+ taskreg_ctx = ctx->outer->outer;
}
- else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP)
- taskreg_ctx = ctx->outer->outer;
- }
- else if (is_taskreg_ctx (ctx->outer))
- taskreg_ctx = ctx->outer;
- if (taskreg_ctx)
- {
- int i;
- tree taskreg_clauses
- = gimple_omp_taskreg_clauses (taskreg_ctx->stmt);
- tree innerc = omp_find_clause (taskreg_clauses,
- OMP_CLAUSE__LOOPTEMP_);
- gcc_assert (innerc);
- for (i = 0; i < fd->collapse; i++)
+ else
{
- innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc),
- OMP_CLAUSE__LOOPTEMP_);
- gcc_assert (innerc);
+ struct omp_for_data outer_fd;
+ omp_extract_for_data (gfor, &outer_fd, NULL);
+ n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
}
+ }
+ else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP)
+ taskreg_ctx = ctx->outer->outer;
+ }
+ else if (is_taskreg_ctx (ctx->outer))
+ taskreg_ctx = ctx->outer;
+ if (taskreg_ctx)
+ {
+ int i;
+ tree taskreg_clauses
+ = gimple_omp_taskreg_clauses (taskreg_ctx->stmt);
+ tree innerc = omp_find_clause (taskreg_clauses,
+ OMP_CLAUSE__LOOPTEMP_);
+ gcc_assert (innerc);
+ int count = fd->collapse;
+ if (fd->non_rect
+ && fd->last_nonrect == fd->first_nonrect + 1)
+ if (tree v = gimple_omp_for_index (fd->for_stmt, fd->last_nonrect))
+ if (!TYPE_UNSIGNED (TREE_TYPE (v)))
+ count += 4;
+ for (i = 0; i < count; i++)
+ {
innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc),
OMP_CLAUSE__LOOPTEMP_);
- if (innerc)
- n2 = fold_convert (TREE_TYPE (n2),
- lookup_decl (OMP_CLAUSE_DECL (innerc),
- taskreg_ctx));
+ gcc_assert (innerc);
}
+ innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ if (innerc)
+ n2 = fold_convert (TREE_TYPE (n2),
+ lookup_decl (OMP_CLAUSE_DECL (innerc),
+ taskreg_ctx));
}
- cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2);
}
+ cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2);
clauses = gimple_omp_for_clauses (fd->for_stmt);
stmts = NULL;
WALK_SUBSTMTS;
case GIMPLE_OMP_FOR:
- if ((gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD)
+ if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_SIMD
&& gimple_omp_for_combined_into_p (stmt))
*handled_ops_p = false;
break;
gimple_seq_add_stmt (body_p, g);
tree cplx = create_tmp_var (build_complex_type (unsigned_type_node, false));
- DECL_GIMPLE_REG_P (cplx) = 1;
g = gimple_build_call_internal (IFN_MUL_OVERFLOW, 2, thread_nump1, twok);
gimple_call_set_lhs (g, cplx);
gimple_seq_add_stmt (body_p, g);
if (fd.collapse > 1
&& TREE_CODE (fd.loop.n2) != INTEGER_CST)
count += fd.collapse - 1;
+ size_t count2 = 0;
+ tree type2 = NULL_TREE;
bool taskreg_for
= (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR
|| gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP);
tree outerc = NULL, *pc = gimple_omp_for_clauses_ptr (stmt);
tree simtc = NULL;
tree clauses = *pc;
+ if (fd.collapse > 1
+ && fd.non_rect
+ && fd.last_nonrect == fd.first_nonrect + 1
+ && TREE_CODE (fd.loop.n2) != INTEGER_CST)
+ if (tree v = gimple_omp_for_index (stmt, fd.last_nonrect))
+ if (!TYPE_UNSIGNED (TREE_TYPE (v)))
+ {
+ v = gimple_omp_for_index (stmt, fd.first_nonrect);
+ type2 = TREE_TYPE (v);
+ count++;
+ count2 = 3;
+ }
if (taskreg_for)
outerc
= omp_find_clause (gimple_omp_taskreg_clauses (ctx->outer->stmt),
if (ctx->simt_stmt)
simtc = omp_find_clause (gimple_omp_for_clauses (ctx->simt_stmt),
OMP_CLAUSE__LOOPTEMP_);
- for (i = 0; i < count; i++)
+ for (i = 0; i < count + count2; i++)
{
tree temp;
if (taskreg_for)
if (ctx->simt_stmt)
temp = OMP_CLAUSE_DECL (simtc);
else
- temp = create_tmp_var (type);
+ temp = create_tmp_var (i >= count ? type2 : type);
insert_decl_map (&ctx->outer->cb, temp, temp);
}
*pc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__LOOPTEMP_);
for (i = 0; i < gimple_omp_for_collapse (stmt); i++)
{
rhs_p = gimple_omp_for_initial_ptr (stmt, i);
- if (!is_gimple_min_invariant (*rhs_p))
+ if (TREE_CODE (*rhs_p) == TREE_VEC)
+ {
+ if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 1)))
+ TREE_VEC_ELT (*rhs_p, 1)
+ = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 1), &cnt_list);
+ if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 2)))
+ TREE_VEC_ELT (*rhs_p, 2)
+ = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 2), &cnt_list);
+ }
+ else if (!is_gimple_min_invariant (*rhs_p))
*rhs_p = get_formal_tmp_var (*rhs_p, &cnt_list);
else if (TREE_CODE (*rhs_p) == ADDR_EXPR)
recompute_tree_invariant_for_addr_expr (*rhs_p);
rhs_p = gimple_omp_for_final_ptr (stmt, i);
- if (!is_gimple_min_invariant (*rhs_p))
+ if (TREE_CODE (*rhs_p) == TREE_VEC)
+ {
+ if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 1)))
+ TREE_VEC_ELT (*rhs_p, 1)
+ = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 1), &cnt_list);
+ if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 2)))
+ TREE_VEC_ELT (*rhs_p, 2)
+ = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 2), &cnt_list);
+ }
+ else if (!is_gimple_min_invariant (*rhs_p))
*rhs_p = get_formal_tmp_var (*rhs_p, &cnt_list);
else if (TREE_CODE (*rhs_p) == ADDR_EXPR)
recompute_tree_invariant_for_addr_expr (*rhs_p);
ctx);
}
- bool phony_loop = (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_GRID_LOOP
- && gimple_omp_for_grid_phony (stmt));
if ((ctx->scan_inclusive || ctx->scan_exclusive)
&& gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR)
- {
- gcc_assert (!phony_loop);
- lower_omp_for_scan (&body, &dlist, stmt, &fd, ctx);
- }
+ lower_omp_for_scan (&body, &dlist, stmt, &fd, ctx);
else
{
- if (!phony_loop)
- gimple_seq_add_stmt (&body, stmt);
+ gimple_seq_add_stmt (&body, stmt);
gimple_seq_add_seq (&body, gimple_omp_body (stmt));
}
- if (!phony_loop)
- gimple_seq_add_stmt (&body, gimple_build_omp_continue (fd.loop.v,
- fd.loop.v));
+ gimple_seq_add_stmt (&body, gimple_build_omp_continue (fd.loop.v,
+ fd.loop.v));
/* After the loop, add exit clauses. */
lower_reduction_clauses (gimple_omp_for_clauses (stmt), &body, &clist, ctx);
body = maybe_catch_exception (body);
- if (!phony_loop)
- {
- /* Region exit marker goes at the end of the loop body. */
- gimple *g = gimple_build_omp_return (fd.have_nowait);
- gimple_seq_add_stmt (&body, g);
+ /* Region exit marker goes at the end of the loop body. */
+ gimple *g = gimple_build_omp_return (fd.have_nowait);
+ gimple_seq_add_stmt (&body, g);
- gimple_seq_add_seq (&body, tred_dlist);
+ gimple_seq_add_seq (&body, tred_dlist);
- maybe_add_implicit_barrier_cancel (ctx, g, &body);
+ maybe_add_implicit_barrier_cancel (ctx, g, &body);
- if (rclauses)
- OMP_CLAUSE_DECL (rclauses) = rtmp;
- }
+ if (rclauses)
+ OMP_CLAUSE_DECL (rclauses) = rtmp;
/* Add OpenACC joining and reduction markers just after the loop. */
if (oacc_tail)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FIRSTPRIVATE)
t = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, src);
else
- t = lang_hooks.decls.omp_clause_copy_ctor (c, dst, src);
+ {
+ if (ctx->allocate_map)
+ if (tree *allocatorp = ctx->allocate_map->get (decl))
+ {
+ tree allocator = *allocatorp;
+ if (TREE_CODE (allocator) != INTEGER_CST)
+ {
+ n = splay_tree_lookup (ctx->sfield_map,
+ (splay_tree_key) allocator);
+ allocator = (tree) n->value;
+ if (tcctx.cb.decl_map)
+ allocator = *tcctx.cb.decl_map->get (allocator);
+ tree a = build_simple_mem_ref_loc (loc, sarg);
+ allocator = omp_build_component_ref (a, allocator);
+ }
+ allocator = fold_convert (pointer_sized_int_node, allocator);
+ tree a = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC);
+ tree align = build_int_cst (size_type_node,
+ DECL_ALIGN_UNIT (decl));
+ tree sz = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (dst)));
+ tree ptr = build_call_expr_loc (loc, a, 3, align, sz,
+ allocator);
+ ptr = fold_convert (TREE_TYPE (dst), ptr);
+ t = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, ptr);
+ append_to_statement_list (t, &list);
+ dst = build_simple_mem_ref_loc (loc, dst);
+ }
+ t = lang_hooks.decls.omp_clause_copy_ctor (c, dst, src);
+ }
append_to_statement_list (t, &list);
break;
case OMP_CLAUSE_PRIVATE:
OMP_CLAUSE_DECL (c) = build_fold_addr_expr (array);
OMP_CLAUSE_CHAIN (c) = *pclauses;
*pclauses = c;
- tree clobber = build_constructor (type, NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (type);
g = gimple_build_assign (array, clobber);
gimple_seq_add_stmt (oseq, g);
}
gimple_seq par_olist = NULL;
gimple_seq par_ilist = NULL;
gimple_seq par_rlist = NULL;
- bool phony_construct = gimple_code (stmt) == GIMPLE_OMP_PARALLEL
- && gimple_omp_parallel_grid_phony (as_a <gomp_parallel *> (stmt));
- if (phony_construct && ctx->record_type)
- {
- gcc_checking_assert (!ctx->receiver_decl);
- ctx->receiver_decl = create_tmp_var
- (build_reference_type (ctx->record_type), ".omp_rec");
- }
lower_rec_input_clauses (clauses, &par_ilist, &par_olist, ctx, NULL);
lower_omp (&par_body, ctx);
- if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL)
+ if (gimple_code (stmt) != GIMPLE_OMP_TASK)
lower_reduction_clauses (clauses, &par_rlist, NULL, ctx);
/* Declare all the variables created by mapping and the variables
if (ctx->record_type)
{
- tree clobber = build_constructor (TREE_TYPE (ctx->sender_decl), NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (TREE_TYPE (ctx->sender_decl));
gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
clobber));
}
gimple_seq_add_stmt (&new_body,
gimple_build_omp_continue (integer_zero_node,
integer_zero_node));
- if (!phony_construct)
- {
- gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
- gimple_omp_set_body (stmt, new_body);
- }
+ gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
+ gimple_omp_set_body (stmt, new_body);
if (dep_bind && gimple_bind_block (par_bind) == NULL_TREE)
bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
bind = gimple_build_bind (NULL, NULL, gimple_bind_block (par_bind));
gsi_replace (gsi_p, dep_bind ? dep_bind : bind, true);
gimple_bind_add_seq (bind, ilist);
- if (!phony_construct)
- gimple_bind_add_stmt (bind, stmt);
- else
- gimple_bind_add_seq (bind, new_body);
+ gimple_bind_add_stmt (bind, stmt);
gimple_bind_add_seq (bind, olist);
pop_gimplify_context (NULL);
case GF_OMP_TARGET_KIND_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
data_region = true;
break;
default:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
break;
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
case GOMP_MAP_LINK:
+ case GOMP_MAP_FORCE_DETACH:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
continue;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (stmt))
+ {
+ gcc_assert (maybe_lookup_field (c, ctx));
+ map_cnt++;
+ continue;
+ }
+
if (!maybe_lookup_field (var, ctx))
continue;
- /* Don't remap oacc parallel reduction variables, because the
+ /* Don't remap compute constructs' reduction variables, because the
intermediate result must be local to each gang. */
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
{
gcc_assert (is_gimple_omp_oacc (ctx->stmt));
if (omp_is_reference (new_var)
- && TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE)
+ && (TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE
+ || DECL_BY_REFERENCE (var)))
{
/* Create a local object to hold the instance
value. */
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_oacc_parallel (ctx))
- goto oacc_firstprivate;
+ gcc_checking_assert (offloaded);
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ /* No 'firstprivate' clauses on OpenACC 'kernels'. */
+ gcc_checking_assert (!is_oacc_kernels (ctx));
+ /* Likewise, on OpenACC 'kernels' decomposed parts. */
+ gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx));
+
+ goto oacc_firstprivate;
+ }
map_cnt++;
var = OMP_CLAUSE_DECL (c);
if (!omp_is_reference (var)
break;
case OMP_CLAUSE_PRIVATE:
+ gcc_checking_assert (offloaded);
if (is_gimple_omp_oacc (ctx->stmt))
- break;
+ {
+ /* No 'private' clauses on OpenACC 'kernels'. */
+ gcc_checking_assert (!is_oacc_kernels (ctx));
+ /* Likewise, on OpenACC 'kernels' decomposed parts. */
+ gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx));
+
+ break;
+ }
var = OMP_CLAUSE_DECL (c);
if (is_variable_sized (var))
{
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
map_cnt++;
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
- else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ && !omp_is_reference (var)
+ && !omp_is_allocatable_or_ptr (var)
+ && !lang_hooks.decls.omp_array_data (var, true))
+ || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
{
tree new_var = lookup_decl (var, ctx);
tree type = build_pointer_type (TREE_TYPE (var));
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
- if (!maybe_lookup_field (ovar, ctx))
+ if (!maybe_lookup_field (ovar, ctx)
+ && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
continue;
}
talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
- if (nc)
+
+ if (nc
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (stmt))
+ {
+ var = lookup_decl_in_outer_ctx (ovar, ctx);
+ x = build_sender_ref (c, ctx);
+ gimplify_assign (x, build_fold_addr_expr (var), &ilist);
+ }
+ else if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
}
else
{
- var = build_fold_addr_expr (var);
+ /* While MAP is handled explicitly by the FE,
+ for 'target update', only the identified is passed. */
+ if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO)
+ && (omp_is_allocatable_or_ptr (var)
+ && omp_check_optional_argument (var, false)))
+ var = build_fold_indirect_ref (var);
+ else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO)
+ || (!omp_is_allocatable_or_ptr (var)
+ && !omp_check_optional_argument (var, false)))
+ var = build_fold_addr_expr (var);
gimplify_assign (x, var, &ilist);
}
}
{
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
s = TREE_TYPE (ovar);
- if (TREE_CODE (s) == REFERENCE_TYPE)
+ if (TREE_CODE (s) == REFERENCE_TYPE
+ || omp_check_optional_argument (ovar, false))
s = TREE_TYPE (s);
s = TYPE_SIZE_UNIT (s);
}
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_oacc_parallel (ctx))
+ if (is_gimple_omp_oacc (ctx->stmt))
goto oacc_firstprivate_map;
ovar = OMP_CLAUSE_DECL (c);
if (omp_is_reference (ovar))
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
var = lookup_decl_in_outer_ctx (ovar, ctx);
- x = build_sender_ref (ovar, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
- tkind = GOMP_MAP_USE_DEVICE_PTR;
+
+ if (lang_hooks.decls.omp_array_data (ovar, true))
+ {
+ tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+ ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT);
+ x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx);
+ }
+ else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
+ {
+ tkind = GOMP_MAP_USE_DEVICE_PTR;
+ x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx);
+ }
else
- tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ {
+ tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ x = build_sender_ref (ovar, ctx);
+ }
+
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ gcc_assert (tkind == GOMP_MAP_USE_DEVICE_PTR);
+
+ if (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (c))
+ tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT;
+ }
+
type = TREE_TYPE (ovar);
- if (TREE_CODE (type) == ARRAY_TYPE)
+ if (lang_hooks.decls.omp_array_data (ovar, true))
+ var = lang_hooks.decls.omp_array_data (ovar, false);
+ else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ && !omp_is_reference (ovar)
+ && !omp_is_allocatable_or_ptr (ovar))
+ || TREE_CODE (type) == ARRAY_TYPE)
var = build_fold_addr_expr (var);
else
{
- if (omp_is_reference (ovar))
+ if (omp_is_reference (ovar)
+ || omp_check_optional_argument (ovar, false)
+ || omp_is_allocatable_or_ptr (ovar))
{
type = TREE_TYPE (type);
- if (TREE_CODE (type) != ARRAY_TYPE)
+ if (TREE_CODE (type) != ARRAY_TYPE
+ && ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
+ && !omp_is_allocatable_or_ptr (ovar))
+ || (omp_is_reference (ovar)
+ && omp_is_allocatable_or_ptr (ovar))))
var = build_simple_mem_ref (var);
var = fold_convert (TREE_TYPE (x), var);
}
}
- gimplify_assign (x, var, &ilist);
+ tree present;
+ present = omp_check_optional_argument (ovar, true);
+ if (present)
+ {
+ tree null_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree notnull_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree opt_arg_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree new_x = unshare_expr (x);
+ gimplify_expr (&present, &ilist, NULL, is_gimple_val,
+ fb_rvalue);
+ gcond *cond = gimple_build_cond_from_tree (present,
+ notnull_label,
+ null_label);
+ gimple_seq_add_stmt (&ilist, cond);
+ gimple_seq_add_stmt (&ilist, gimple_build_label (null_label));
+ gimplify_assign (new_x, null_pointer_node, &ilist);
+ gimple_seq_add_stmt (&ilist, gimple_build_goto (opt_arg_label));
+ gimple_seq_add_stmt (&ilist,
+ gimple_build_label (notnull_label));
+ gimplify_assign (x, var, &ilist);
+ gimple_seq_add_stmt (&ilist,
+ gimple_build_label (opt_arg_label));
+ }
+ else
+ gimplify_assign (x, var, &ilist);
s = size_int (0);
purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
&initlist, true, NULL_TREE);
gimple_seq_add_seq (&ilist, initlist);
- tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, i)),
- NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (TREE_TYPE (TREE_VEC_ELT (t, i)));
gimple_seq_add_stmt (&olist,
gimple_build_assign (TREE_VEC_ELT (t, i),
clobber));
}
- tree clobber = build_constructor (ctx->record_type, NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
+ tree clobber = build_clobber (ctx->record_type);
gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
clobber));
}
}
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
+ tree new_var;
+ gimple_seq assign_body;
+ bool is_array_data;
+ bool do_optional_check;
+ assign_body = NULL;
+ do_optional_check = false;
var = OMP_CLAUSE_DECL (c);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
- x = build_sender_ref (var, ctx);
+ is_array_data = lang_hooks.decls.omp_array_data (var, true) != NULL;
+
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
+ x = build_sender_ref (is_array_data
+ ? (splay_tree_key) &DECL_NAME (var)
+ : (splay_tree_key) &DECL_UID (var), ctx);
else
x = build_receiver_ref (var, false, ctx);
- if (is_variable_sized (var))
+
+ if (is_array_data)
+ {
+ bool is_ref = omp_is_reference (var);
+ do_optional_check = true;
+ /* First, we copy the descriptor data from the host; then
+ we update its data to point to the target address. */
+ new_var = lookup_decl (var, ctx);
+ new_var = DECL_VALUE_EXPR (new_var);
+ tree v = new_var;
+
+ if (is_ref)
+ {
+ var = build_fold_indirect_ref (var);
+ gimplify_expr (&var, &assign_body, NULL, is_gimple_val,
+ fb_rvalue);
+ v = create_tmp_var_raw (TREE_TYPE (var), get_name (var));
+ gimple_add_tmp_var (v);
+ TREE_ADDRESSABLE (v) = 1;
+ gimple_seq_add_stmt (&assign_body,
+ gimple_build_assign (v, var));
+ tree rhs = build_fold_addr_expr (v);
+ gimple_seq_add_stmt (&assign_body,
+ gimple_build_assign (new_var, rhs));
+ }
+ else
+ gimple_seq_add_stmt (&assign_body,
+ gimple_build_assign (new_var, var));
+
+ tree v2 = lang_hooks.decls.omp_array_data (unshare_expr (v), false);
+ gcc_assert (v2);
+ gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&assign_body,
+ gimple_build_assign (v2, x));
+ }
+ else if (is_variable_sized (var))
{
tree pvar = DECL_VALUE_EXPR (var);
gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
pvar = TREE_OPERAND (pvar, 0);
gcc_assert (DECL_P (pvar));
- tree new_var = lookup_decl (pvar, ctx);
- gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
- gimple_seq_add_stmt (&new_body,
+ new_var = lookup_decl (pvar, ctx);
+ gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&assign_body,
gimple_build_assign (new_var, x));
}
- else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ && !omp_is_reference (var)
+ && !omp_is_allocatable_or_ptr (var))
+ || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
{
- tree new_var = lookup_decl (var, ctx);
+ new_var = lookup_decl (var, ctx);
new_var = DECL_VALUE_EXPR (new_var);
gcc_assert (TREE_CODE (new_var) == MEM_REF);
new_var = TREE_OPERAND (new_var, 0);
gcc_assert (DECL_P (new_var));
- gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
- gimple_seq_add_stmt (&new_body,
+ gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&assign_body,
gimple_build_assign (new_var, x));
}
else
{
tree type = TREE_TYPE (var);
- tree new_var = lookup_decl (var, ctx);
+ new_var = lookup_decl (var, ctx);
if (omp_is_reference (var))
{
type = TREE_TYPE (type);
- if (TREE_CODE (type) != ARRAY_TYPE)
+ if (TREE_CODE (type) != ARRAY_TYPE
+ && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
+ || (omp_is_reference (var)
+ && omp_is_allocatable_or_ptr (var))))
{
tree v = create_tmp_var_raw (type, get_name (var));
gimple_add_tmp_var (v);
TREE_ADDRESSABLE (v) = 1;
x = fold_convert (type, x);
- gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+ gimplify_expr (&x, &assign_body, NULL, is_gimple_val,
fb_rvalue);
- gimple_seq_add_stmt (&new_body,
+ gimple_seq_add_stmt (&assign_body,
gimple_build_assign (v, x));
x = build_fold_addr_expr (v);
+ do_optional_check = true;
}
}
new_var = DECL_VALUE_EXPR (new_var);
x = fold_convert (TREE_TYPE (new_var), x);
- gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
- gimple_seq_add_stmt (&new_body,
+ gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&assign_body,
gimple_build_assign (new_var, x));
}
+ tree present;
+ present = (do_optional_check
+ ? omp_check_optional_argument (OMP_CLAUSE_DECL (c), true)
+ : NULL_TREE);
+ if (present)
+ {
+ tree null_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree notnull_label = create_artificial_label (UNKNOWN_LOCATION);
+ tree opt_arg_label = create_artificial_label (UNKNOWN_LOCATION);
+ glabel *null_glabel = gimple_build_label (null_label);
+ glabel *notnull_glabel = gimple_build_label (notnull_label);
+ ggoto *opt_arg_ggoto = gimple_build_goto (opt_arg_label);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+ fb_rvalue);
+ gimplify_expr (&present, &new_body, NULL, is_gimple_val,
+ fb_rvalue);
+ gcond *cond = gimple_build_cond_from_tree (present,
+ notnull_label,
+ null_label);
+ gimple_seq_add_stmt (&new_body, cond);
+ gimple_seq_add_stmt (&new_body, null_glabel);
+ gimplify_assign (new_var, null_pointer_node, &new_body);
+ gimple_seq_add_stmt (&new_body, opt_arg_ggoto);
+ gimple_seq_add_stmt (&new_body, notnull_glabel);
+ gimple_seq_add_seq (&new_body, assign_body);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_label (opt_arg_label));
+ }
+ else
+ gimple_seq_add_seq (&new_body, assign_body);
break;
}
/* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
gimple_seq fork_seq = NULL;
gimple_seq join_seq = NULL;
- if (is_oacc_parallel (ctx))
+ if (offloaded && is_gimple_omp_oacc (ctx->stmt))
{
/* If there are reductions on the offloaded region itself, treat
them as a dummy GANG loop. */
lower_omp (gimple_omp_body_ptr (teams_stmt), ctx);
lower_reduction_clauses (gimple_omp_teams_clauses (teams_stmt), &olist,
NULL, ctx);
- if (!gimple_omp_teams_grid_phony (teams_stmt))
- {
- gimple_seq_add_stmt (&bind_body, teams_stmt);
- location_t loc = gimple_location (teams_stmt);
- tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS);
- gimple *call = gimple_build_call (decl, 2, num_teams, thread_limit);
- gimple_set_location (call, loc);
- gimple_seq_add_stmt (&bind_body, call);
- }
+ gimple_seq_add_stmt (&bind_body, teams_stmt);
+
+ location_t loc = gimple_location (teams_stmt);
+ tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS);
+ gimple *call = gimple_build_call (decl, 2, num_teams, thread_limit);
+ gimple_set_location (call, loc);
+ gimple_seq_add_stmt (&bind_body, call);
gimple_seq_add_seq (&bind_body, gimple_omp_body (teams_stmt));
gimple_omp_set_body (teams_stmt, NULL);
gimple_seq_add_seq (&bind_body, olist);
gimple_seq_add_seq (&bind_body, dlist);
- if (!gimple_omp_teams_grid_phony (teams_stmt))
- gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true));
+ gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true));
gimple_bind_set_body (bind, bind_body);
pop_gimplify_context (bind);
TREE_USED (block) = 1;
}
-/* Expand code within an artificial GIMPLE_OMP_GRID_BODY OMP construct. */
-
-static void
-lower_omp_grid_body (gimple_stmt_iterator *gsi_p, omp_context *ctx)
-{
- gimple *stmt = gsi_stmt (*gsi_p);
- lower_omp (gimple_omp_body_ptr (stmt), ctx);
- gimple_seq_add_stmt (gimple_omp_body_ptr (stmt),
- gimple_build_omp_return (false));
-}
-
-
/* 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 OMP context, but with task_shared_vars set. */
else
lower_omp_teams (gsi_p, ctx);
break;
- case GIMPLE_OMP_GRID_BODY:
- ctx = maybe_lookup_ctx (stmt);
- gcc_assert (ctx);
- lower_omp_grid_body (gsi_p, ctx);
- break;
case GIMPLE_CALL:
tree fndecl;
call_stmt = as_a <gcall *> (stmt);
{
tree clauses;
if (up->combined_into_simd_safelen1)
- up = up->outer;
+ {
+ up = up->outer;
+ if (gimple_code (up->stmt) == GIMPLE_OMP_SCAN)
+ up = up->outer;
+ }
if (gimple_code (up->stmt) == GIMPLE_OMP_FOR)
clauses = gimple_omp_for_clauses (up->stmt);
else
body = gimple_body (current_function_decl);
- if (hsa_gen_requested_p ())
- omp_grid_gridify_all_targets (&body);
-
scan_omp (&body, NULL);
gcc_assert (taskreg_nesting_level == 0);
FOR_EACH_VEC_ELT (taskreg_contexts, i, ctx)
all_contexts = NULL;
}
BITMAP_FREE (task_shared_vars);
+ BITMAP_FREE (global_nonaddressable_vars);
/* If current function is a method, remove artificial dummy VAR_DECL created
for non-static data member privatization, they aren't needed for
projects / gcc.git / blobdiff