+2016-12-14 Martin Jambor <mjambor@suse.cz>
+
+ * omp-offload.c: Fix coding style.
+ * omp-expand.c: Likewise.
+ * omp-general.c: Likewise.
+ * omp-grid.c: Likewise.
+ * omp-low.c: Fix coding style of parts touched by the
+ previous splitting patch.
+
2016-12-14 Martin Jambor <mjambor@suse.cz>
* omp-general.h: New file.
Is lowered into:
- # BLOCK 2 (PAR_ENTRY_BB)
+ # BLOCK 2 (PAR_ENTRY_BB)
.omp_data_o.i = i;
#pragma omp parallel [child fn: bar.omp_fn.0 ( ..., D.1598)
if (TREE_CODE (arg) == ADDR_EXPR
&& TREE_OPERAND (arg, 0)
- == gimple_omp_taskreg_data_arg (entry_stmt))
+ == gimple_omp_taskreg_data_arg (entry_stmt))
{
parcopy_stmt = stmt;
break;
gsi_remove (&gsi, true);
else
{
- /* ?? Is setting the subcode really necessary ?? */
+ /* ?? Is setting the subcode really necessary ?? */
gimple_omp_set_subcode (parcopy_stmt, TREE_CODE (arg));
gimple_assign_set_rhs1 (parcopy_stmt, arg);
}
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. */
+ loop tree arrange for the child function to fixup loops. */
if (loops_state_satisfies_p (LOOPS_NEED_FIXUP))
child_cfun->x_current_loops->state |= LOOPS_NEED_FIXUP;
struct oacc_collapse
{
- tree base; /* Base value. */
+ tree base; /* Base value. */
tree iters; /* Number of steps. */
tree step; /* step size. */
};
e = force_gimple_operand_gsi (gsi, e, true, NULL_TREE,
true, GSI_SAME_STMT);
- /* Convert the step, avoiding possible unsigned->signed overflow. */
+ /* Convert the step, avoiding possible unsigned->signed overflow. */
negating = !up && TYPE_UNSIGNED (TREE_TYPE (s));
if (negating)
s = fold_build1 (NEGATE_EXPR, TREE_TYPE (s), s);
s = force_gimple_operand_gsi (gsi, s, true, NULL_TREE,
true, GSI_SAME_STMT);
- /* Determine the range, avoiding possible unsigned->signed overflow. */
+ /* Determine the range, avoiding possible unsigned->signed overflow. */
negating = !up && TYPE_UNSIGNED (iter_type);
expr = fold_build2 (MINUS_EXPR, plus_type,
fold_convert (plus_type, negating ? b : e),
t = fold_build2 (NE_EXPR, boolean_type_node,
t, build_int_cst (TREE_TYPE (t), 0));
t = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
- true, GSI_SAME_STMT);
+ true, GSI_SAME_STMT);
if (arr && !TREE_STATIC (arr))
{
tree clobber = build_constructor (TREE_TYPE (arr), NULL);
if the loop is not entered
L0:
s0 = (trip * nthreads + threadid) * CHUNK;
- e0 = min(s0 + CHUNK, n);
+ e0 = min (s0 + CHUNK, n);
if (s0 < n) goto L1; else goto L4;
L1:
V = s0 * STEP + N1;
find_edge (cont_bb, trip_update_bb)->flags
= se ? EDGE_FALSE_VALUE : EDGE_FALLTHRU;
- redirect_edge_and_branch (single_succ_edge (trip_update_bb), iter_part_bb);
+ redirect_edge_and_branch (single_succ_edge (trip_update_bb),
+ iter_part_bb);
}
if (gimple_in_ssa_p (cfun))
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)));
+ gomp_parallel *par_stmt
+ = as_a <gomp_parallel *> (last_stmt (region->outer->entry));
+ tree child_fndecl = gimple_omp_parallel_child_fn (par_stmt);
tree t, low_val = NULL_TREE, high_val = NULL_TREE;
for (t = DECL_ARGUMENTS (child_fndecl); t; t = TREE_CHAIN (t))
{
the loop. */
if ((flag_tree_loop_vectorize
|| (!global_options_set.x_flag_tree_loop_vectorize
- && !global_options_set.x_flag_tree_vectorize))
+ && !global_options_set.x_flag_tree_vectorize))
&& flag_tree_loop_optimize
&& loop->safelen > 1)
{
b = force_gimple_operand_gsi (&gsi, b, true, NULL_TREE, true, GSI_SAME_STMT);
e = force_gimple_operand_gsi (&gsi, e, true, NULL_TREE, true, GSI_SAME_STMT);
- /* Convert the step, avoiding possible unsigned->signed overflow. */
+ /* Convert the step, avoiding possible unsigned->signed overflow. */
negating = !up && TYPE_UNSIGNED (TREE_TYPE (s));
if (negating)
s = fold_build1 (NEGATE_EXPR, TREE_TYPE (s), s);
expr = fold_convert (diff_type, chunk_size);
chunk_size = force_gimple_operand_gsi (&gsi, expr, true,
NULL_TREE, true, GSI_SAME_STMT);
- /* Determine the range, avoiding possible unsigned->signed overflow. */
+ /* Determine the range, avoiding possible unsigned->signed overflow. */
negating = !up && TYPE_UNSIGNED (iter_type);
expr = fold_build2 (MINUS_EXPR, plus_type,
fold_convert (plus_type, negating ? b : e),
/* Remove the GIMPLE_OMP_FOR. */
gsi_remove (&gsi, true);
- /* Fixup edges from head_bb */
+ /* Fixup edges from head_bb. */
be = BRANCH_EDGE (head_bb);
fte = FALLTHRU_EDGE (head_bb);
be->flags |= EDGE_FALSE_VALUE;
/* Remove the GIMPLE_OMP_CONTINUE. */
gsi_remove (&gsi, true);
- /* Fixup edges from cont_bb */
+ /* Fixup edges from cont_bb. */
be = BRANCH_EDGE (cont_bb);
fte = FALLTHRU_EDGE (cont_bb);
be->flags |= EDGE_TRUE_VALUE;
{
/* Split the beginning of exit_bb to make bottom_bb. We
need to insert a nop at the start, because splitting is
- after a stmt, not before. */
+ after a stmt, not before. */
gsi = gsi_start_bb (exit_bb);
stmt = gimple_build_nop ();
gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
gsi_insert_after (&gsi, gimple_build_cond_empty (expr),
GSI_CONTINUE_LINKING);
- /* Fixup edges from bottom_bb. */
+ /* Fixup edges from bottom_bb. */
split->flags ^= EDGE_FALLTHRU | EDGE_FALSE_VALUE;
make_edge (bottom_bb, head_bb, EDGE_TRUE_VALUE);
}
gsi_insert_before (&gsi, ass, GSI_SAME_STMT);
}
- /* Remove the OMP_RETURN. */
+ /* Remove the OMP_RETURN. */
gsi_remove (&gsi, true);
if (cont_bb)
si = gsi_last_bb (e->dest);
l2 = NULL_TREE;
if (gsi_end_p (si)
- || gimple_code (gsi_stmt (si)) != GIMPLE_OMP_SECTION)
+ || gimple_code (gsi_stmt (si)) != GIMPLE_OMP_SECTION)
l2 = gimple_block_label (e->dest);
else
FOR_EACH_EDGE (e, ei, l0_bb->succs)
oldval = *addr;
repeat:
- newval = rhs; // with oldval replacing *addr in rhs
+ newval = rhs; // with oldval replacing *addr in rhs
oldval = __sync_val_compare_and_swap (addr, oldval, newval);
if (oldval != newval)
goto repeat;
if (iaddr == addr)
storedi = stored_val;
else
- storedi =
- force_gimple_operand_gsi (&si,
- build1 (VIEW_CONVERT_EXPR, itype,
- stored_val), true, NULL_TREE, true,
- GSI_SAME_STMT);
+ storedi
+ = force_gimple_operand_gsi (&si,
+ build1 (VIEW_CONVERT_EXPR, itype,
+ stored_val), true, NULL_TREE, true,
+ GSI_SAME_STMT);
/* Build the compare&swap statement. */
new_storedi = build_call_expr (cmpxchg, 3, iaddr, loadedi, storedi);
/* Note that we always perform the comparison as an integer, even for
floating point. This allows the atomic operation to properly
succeed even with NaNs and -0.0. */
- stmt = gimple_build_cond_empty
- (build2 (NE_EXPR, boolean_type_node,
- new_storedi, old_vali));
+ tree ne = build2 (NE_EXPR, boolean_type_node, new_storedi, old_vali);
+ stmt = gimple_build_cond_empty (ne);
gsi_insert_before (&si, stmt, GSI_SAME_STMT);
/* Update cfg. */
/* A subroutine of expand_omp_atomic. Implement the atomic operation as:
- GOMP_atomic_start ();
- *addr = rhs;
- GOMP_atomic_end ();
+ GOMP_atomic_start ();
+ *addr = rhs;
+ GOMP_atomic_end ();
The result is not globally atomic, but works so long as all parallel
references are within #pragma omp atomic directives. According to
}
/* Expand an GIMPLE_OMP_ATOMIC statement. We try to expand
- using expand_omp_atomic_fetch_op. If it failed, we try to
+ using expand_omp_atomic_fetch_op. If it failed, we try to
call expand_omp_atomic_pipeline, and if it fails too, the
ultimate fallback is wrapping the operation in a mutex
(expand_omp_atomic_mutex). REGION is the atomic region built
if (nr_outer_loops != 1)
return;
- for (struct loop *loop = single_outer->inner; loop != NULL; loop = loop->inner)
+ for (struct loop *loop = single_outer->inner;
+ loop != NULL;
+ loop = loop->inner)
if (loop->next)
return;
}
}
-/* Create an array of arguments that is then passed to GOMP_target. */
+/* Create an array of arguments that is then passed to GOMP_target. */
static tree
get_target_arguments (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt)
if (omp_find_clause (gimple_omp_target_clauses (tgt_stmt),
OMP_CLAUSE__GRIDDIM_))
{
- t = get_target_argument_identifier (GOMP_DEVICE_HSA, true,
- GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES);
+ int id = GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES;
+ t = get_target_argument_identifier (GOMP_DEVICE_HSA, true, id);
args.quick_push (t);
args.quick_push (grid_get_kernel_launch_attributes (gsi, tgt_stmt));
}
/* Expand KFOR loop as a HSA grifidied kernel, i.e. as a body only with
iteration variable derived from the thread number. INTRA_GROUP means this
is an expansion of a loop iterating over work-items within a separate
- iteration over groups. */
+ iteration over groups. */
static void
grid_expand_omp_for_loop (struct omp_region *kfor, bool intra_group)
size_t collapse = gimple_omp_for_collapse (for_stmt);
struct omp_for_data_loop *loops
= XALLOCAVEC (struct omp_for_data_loop,
- gimple_omp_for_collapse (for_stmt));
+ gimple_omp_for_collapse (for_stmt));
struct omp_for_data fd;
remove_edge (BRANCH_EDGE (kfor->entry));
gassign *assign_stmt = gimple_build_assign (startvar, t);
gsi_insert_before (&gsi, assign_stmt, GSI_SAME_STMT);
}
- /* Remove the omp for statement */
+ /* Remove the omp for statement. */
gsi = gsi_last_bb (kfor->entry);
gsi_remove (&gsi, true);
}
/* If TARGET region contains a kernel body for loop, remove its region from the
- TARGET and expand it in HSA gridified kernel fashion. */
+ TARGET and expand it in HSA gridified kernel fashion. */
static void
grid_expand_target_grid_body (struct omp_region *target)
gcc_assert (omp_find_clause (gimple_omp_target_clauses (tgt_stmt),
OMP_CLAUSE__GRIDDIM_));
- tree inside_block = gimple_block (first_stmt (single_succ (gpukernel->entry)));
+ tree inside_block
+ = gimple_block (first_stmt (single_succ (gpukernel->entry)));
*pp = gpukernel->next;
for (pp = &gpukernel->inner; *pp; pp = &(*pp)->next)
if ((*pp)->type == GIMPLE_OMP_FOR)
grid_expand_omp_for_loop (kfor, false);
- /* Remove the omp for statement */
+ /* Remove the omp for statement. */
gimple_stmt_iterator gsi = gsi_last_bb (gpukernel->entry);
gsi_remove (&gsi, true);
/* Replace the GIMPLE_OMP_RETURN at the end of the kernel region with a real
gimple *inner_stmt = NULL;
/* First, determine whether this is a combined parallel+workshare
- region. */
+ region. */
if (region->type == GIMPLE_OMP_PARALLEL)
determine_parallel_type (region);
else if (region->type == GIMPLE_OMP_TARGET)
|| !flag_tree_loop_optimize
|| (!flag_tree_loop_vectorize
&& (global_options_set.x_flag_tree_loop_vectorize
- || global_options_set.x_flag_tree_vectorize)))
+ || global_options_set.x_flag_tree_vectorize)))
return 1;
int vf = 1;
if (!optimize)
return 0;
if (ENABLE_OFFLOADING)
- for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c; )
+ for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
{
if (!strncmp (c, "nvptx", strlen ("nvptx")))
return 32;
represented as a list of INTEGER_CST. Those that are runtime
exprs are represented as an INTEGER_CST of zero.
- TOOO. Normally the attribute will just contain a single such list. If
+ TODO: Normally the attribute will just contain a single such list. If
however it contains a list of lists, this will represent the use of
device_type. Each member of the outer list is an assoc list of
dimensions, keyed by the device type. The first entry will be the
oacc_build_routine_dims (tree clauses)
{
/* Must match GOMP_DIM ordering. */
- static const omp_clause_code ids[] =
- {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
+ static const omp_clause_code ids[]
+ = {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
int ix;
int level = -1;
/* When dealing with a gridified loop, we need to check up to three collapsed
iteration variables but they are not actually captured in this fd.
Fortunately, we can easily rely on HSA builtins to get this
- information. */
+ information. */
tree id, size;
if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_GRID_LOOP
}
/* Structure describing the basic properties of the loop we ara analyzing
- whether it can be gridified and when it is gridified. */
+ whether it can be gridified and when it is gridified. */
struct grid_prop
{
continue;
if (gbind *bind = dyn_cast <gbind *> (stmt))
{
- if (!grid_find_single_omp_among_assignments_1 (gimple_bind_body (bind),
- grid, name, ret))
+ gimple_seq bind_body = gimple_bind_body (bind);
+ if (!grid_find_single_omp_among_assignments_1 (bind_body, grid, name,
+ ret))
return false;
}
else if (is_gimple_omp (stmt))
/* Examine clauses and the body of omp loop statement GFOR and if something
prevents gridification, issue a missed-optimization diagnostics and return
- false, otherwise return true. GRID describes hitherto discovered properties
+ false, otherwise return true. GRID describes hitherto discovered properties
of the loop that is evaluated for possible gridification. */
static bool
/* Given distribute omp construct represented by DIST, which in the original
source forms a compound construct with a looping construct, return true if it
- can be turned into a gridified HSA kernel. Otherwise return false. GRID
+ can be turned into a gridified HSA kernel. Otherwise return false. GRID
describes hitherto discovered properties of the loop that is evaluated for
possible gridification. */
/* Given an omp loop statement GFOR, return true if it can participate in
tiling gridification, i.e. in one where the distribute and parallel for
loops do not form a compound statement. GRID describes hitherto discovered
- properties of the loop that is evaluated for possible gridification. */
+ properties of the loop that is evaluated for possible gridification. */
static bool
grid_gfor_follows_tiling_pattern (gomp_for *gfor, grid_prop *grid)
/* Given a sequence of statements within a distribute omp construct or a
parallel construct, which in the original source does not form a compound
construct with a looping construct, return true if it does not prevent us
- from turning it into a gridified HSA kernel. Otherwise return false. GRID
+ from turning it into a gridified HSA kernel. Otherwise return false. GRID
describes hitherto discovered properties of the loop that is evaluated for
possible gridification. IN_PARALLEL must be true if seq is within a
parallel construct and flase if it is only within a distribute
their uses. Fortunately, we do not have to do this because if they are
not addressable, it means they are not used in atomic or parallel
statements and so relaxed GPU consistency rules mean we can just keep them
- private. */
+ private. */
if (!TREE_ADDRESSABLE (var))
return;
(gimple_bind_body (bind), dst, tgt_bind, var_segment, wi);
if (var_segment != GRID_SEGMENT_PRIVATE)
- for (tree var = gimple_bind_vars (bind); var; var = DECL_CHAIN (var))
+ for (tree var = gimple_bind_vars (bind);
+ var;
+ var = DECL_CHAIN (var))
grid_mark_variable_segment (var, var_segment);
if (r)
return r;
gcc_assert (teams);
gimple_omp_teams_set_grid_phony (teams, true);
stmt = grid_copy_leading_local_assignments (gimple_omp_body (teams), dst,
- tgt_bind, GRID_SEGMENT_GLOBAL, wi);
+ tgt_bind, GRID_SEGMENT_GLOBAL,
+ wi);
gcc_checking_assert (stmt);
gomp_for *dist = dyn_cast <gomp_for *> (stmt);
gcc_assert (dist);
gomp_for *inner_loop = grid_process_kernel_body_copy (&grid, kernel_seq, gsi,
tgt_bind, &wi);
- gbind *old_bind = as_a <gbind *> (gimple_seq_first (gimple_omp_body (target)));
+ gbind *old_bind
+ = as_a <gbind *> (gimple_seq_first (gimple_omp_body (target)));
gbind *new_bind = as_a <gbind *> (gimple_seq_first (kernel_seq));
tree new_block = gimple_bind_block (new_bind);
tree enc_block = BLOCK_SUPERCONTEXT (gimple_bind_block (old_bind));
else
t = fold_build2 (TRUNC_DIV_EXPR, itype, t, step);
if (grid.tiling)
- {
- if (cond_code == GT_EXPR)
- step = fold_build1 (NEGATE_EXPR, itype, step);
- t = fold_build2 (MULT_EXPR, itype, t, step);
- }
+ {
+ if (cond_code == GT_EXPR)
+ step = fold_build1 (NEGATE_EXPR, itype, step);
+ t = fold_build2 (MULT_EXPR, itype, t, step);
+ }
tree gs = fold_convert (uint32_type_node, t);
gimple_seq tmpseq = NULL;
return;
}
-/* Walker function doing all the work for create_target_kernels. */
+/* Walker function doing all the work for create_target_kernels. */
static tree
grid_gridify_all_targets_stmt (gimple_stmt_iterator *gsi,
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
{
bool ok = false;
-
+
if (ctx)
switch (gimple_code (ctx->stmt))
{
&& GET_MODE_SIZE (vmode) < vs
&& GET_MODE_2XWIDER_MODE (vmode) != VOIDmode)
vmode = GET_MODE_2XWIDER_MODE (vmode);
-
+
tree type = lang_hooks.types.type_for_mode (mode, 1);
if (type == NULL_TREE || TYPE_MODE (type) != mode)
continue;
var = orig;
incoming = outgoing = var;
-
+
if (!inner)
{
/* See if an outer construct also reduces this variable. */
default:
goto do_lookup;
}
-
+
outer = probe;
for (; cls; cls = OMP_CLAUSE_CHAIN (cls))
if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_REDUCTION
}
incoming = outgoing = (t ? t : orig);
}
-
+
has_outer_reduction:;
}
if (!ref_to_res)
ref_to_res = integer_zero_node;
- if (omp_is_reference (orig))
+ if (omp_is_reference (orig))
{
tree type = TREE_TYPE (var);
const char *id = IDENTIFIER_POINTER (DECL_NAME (var));
new_body = maybe_catch_exception (new_body);
- t = gimple_build_omp_return
- (!!omp_find_clause (gimple_omp_sections_clauses (stmt),
- OMP_CLAUSE_NOWAIT));
+ bool nowait = omp_find_clause (gimple_omp_sections_clauses (stmt),
+ OMP_CLAUSE_NOWAIT) != NULL_TREE;
+ t = gimple_build_omp_return (nowait);
gimple_seq_add_stmt (&new_body, t);
maybe_add_implicit_barrier_cancel (ctx, &new_body);
lower_omp_single (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree block;
- 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;
bind_body = maybe_catch_exception (bind_body);
- t = gimple_build_omp_return
- (!!omp_find_clause (gimple_omp_single_clauses (single_stmt),
- OMP_CLAUSE_NOWAIT));
- gimple_seq_add_stmt (&bind_body_tail, t);
+ bool nowait = omp_find_clause (gimple_omp_single_clauses (single_stmt),
+ OMP_CLAUSE_NOWAIT) != NULL_TREE;
+ gimple *g = gimple_build_omp_return (nowait);
+ gimple_seq_add_stmt (&bind_body_tail, g);
maybe_add_implicit_barrier_cancel (ctx, &bind_body_tail);
if (ctx->record_type)
{
}
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));
+ lock = build_call_expr_loc (loc, lock, 1,
+ build_fold_addr_expr_loc (loc, decl));
unlock = builtin_decl_explicit (BUILT_IN_GOMP_CRITICAL_NAME_END);
unlock = build_call_expr_loc (loc, unlock, 1,
gimple_omp_for_clauses (stmt),
&oacc_head, &oacc_tail, ctx);
- /* Add OpenACC partitioning and reduction markers just before the loop */
+ /* Add OpenACC partitioning and reduction markers just before the loop. */
if (oacc_head)
gimple_seq_add_seq (&body, oacc_head);
-
+
lower_omp_for_lastprivate (&fd, &body, &dlist, ctx);
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR)
lower_omp (gimple_try_cleanup_ptr (stmt), ctx);
break;
case GIMPLE_TRANSACTION:
- lower_omp (gimple_transaction_body_ptr (
- as_a <gtransaction *> (stmt)),
+ lower_omp (gimple_transaction_body_ptr (as_a <gtransaction *> (stmt)),
ctx);
break;
case GIMPLE_BIND:
kind = "OpenMP";
}
- /*
- Previously we kept track of the label's entire context in diagnose_sb_[12]
+ /* Previously we kept track of the label's entire context in diagnose_sb_[12]
so we could traverse it and issue a correct "exit" or "enter" error
message upon a structured block violation.
We built the context by building a list with tree_cons'ing, but there is
no easy counterpart in gimple tuples. It seems like far too much work
for issuing exit/enter error messages. If someone really misses the
- distinct error message... patches welcome.
- */
+ distinct error message... patches welcome. */
#if 0
/* Try to avoid confusing the user by producing and error message
gcall *marker; /* Initial head marker. */
- gcall *heads[GOMP_DIM_MAX]; /* Head marker functions. */
- gcall *tails[GOMP_DIM_MAX]; /* Tail marker functions. */
+ gcall *heads[GOMP_DIM_MAX]; /* Head marker functions. */
+ gcall *tails[GOMP_DIM_MAX]; /* Tail marker functions. */
tree routine; /* Pseudo-loop enclosing a routine. */
operate on adjacent iterations. At the worker and gang level,
each gang/warp executes a set of contiguous iterations. Chunking
can override this such that each iteration engine executes a
- contiguous chunk, and then moves on to stride to the next chunk. */
+ contiguous chunk, and then moves on to stride to the next chunk. */
static void
oacc_xform_loop (gcall *call)
/* Validate and update the dimensions for offloaded FN. ATTRS is the
raw attribute. DIMS is an array of dimensions, which is filled in.
LEVEL is the partitioning level of a routine, or -1 for an offload
- region itself. USED is the mask of partitioned execution in the
+ region itself. USED is the mask of partitioned execution in the
function. */
static void
loop->marker = marker;
/* TODO: This is where device_type flattening would occur for the loop
- flags. */
+ flags. */
loop->flags = TREE_INT_CST_LOW (gimple_call_arg (marker, 3));
do
{
if (loop->child)
- loop->child = oacc_loop_sibling_nreverse (loop->child);
+ loop->child = oacc_loop_sibling_nreverse (loop->child);
oacc_loop *next = loop->sibling;
loop->sibling = last;
loop->flags &= ~OLF_AUTO;
if (seq_par)
{
- loop->flags &=
- ~((GOMP_DIM_MASK (GOMP_DIM_MAX) - 1) << OLF_DIM_BASE);
+ loop->flags
+ &= ~((GOMP_DIM_MASK (GOMP_DIM_MAX) - 1) << OLF_DIM_BASE);
this_mask = 0;
}
}
/* Allocate the loop at the innermost available level. */
unsigned this_mask = 0;
- /* Determine the outermost partitioning used within this loop. */
+ /* Determine the outermost partitioning used within this loop. */
this_mask = loop->inner | GOMP_DIM_MASK (GOMP_DIM_MAX);
this_mask = least_bit_hwi (this_mask);
/* Pick the partitioning just inside that one. */
this_mask >>= 1;
- /* And avoid picking one use by an outer loop. */
+ /* And avoid picking one use by an outer loop. */
this_mask &= ~outer_mask;
if (!this_mask && noisy)
return changed;
}
-/* Default dimension bound is unknown on accelerator and 1 on host. */
+/* Default dimension bound is unknown on accelerator and 1 on host. */
int
default_goacc_dim_limit (int ARG_UNUSED (axis))
projects / gcc.git / commitdiff