On Wed, Oct 21, 2015 at 03:49:08PM -0400, Nathan Sidwell wrote:
> This patch is the device-specific half of the previous patch. It processes
> the partition head & tail markers and loop abstraction functions inserted
> during omp lowering.
> > I don't see anything that would e.g. set the various flags that e.g. OpenMP
> > #pragma omp simd or Cilk+ #pragma simd sets, like loop->safelen,
> > loop->force_vectorize, maybe loop->simduid and promote some vars to simduid
> > arrays if that is relevant to OpenACC.
> It won't convert them into such representations.
Can you fix that incrementally? I'd expect that code marked with acc loop
vector
can't have loop carried backward lexical dependencies, at least not within
the adjacent number of iterations specified in vector clause?
> +/* Find the number of threads (POS = false), or thread number (POS =
> + tre) for an OpenACC region partitioned as MASK. Setup code
Typo, tre -> true.
> +static tree
> +oacc_thread_numbers (bool pos, int mask, gimple_seq *seq)
> +{
> + tree res = pos ? NULL_TREE : build_int_cst (unsigned_type_node, 1);
Formatting, too many spaces.
> + if (res == NULL_TREE)
> + res = build_int_cst (integer_type_node, 0);
integer_zero_node ?
> +/* Transform IFN_GOACC_LOOP calls to actual code. See
> + expand_oacc_for for where these are generated. At the vector
> + level, we stride loops, such that each member of a warp will
Too many spaces before member.
> + gimple_stmt_iterator gsi = gsi_for_stmt (call);
> + unsigned code = (unsigned)TREE_INT_CST_LOW (gimple_call_arg (call, 0));
Missing space before T.
> + tree dir = gimple_call_arg (call, 1);
> + tree range = gimple_call_arg (call, 2);
> + tree step = gimple_call_arg (call, 3);
> + tree chunk_size = NULL_TREE;
> + unsigned mask = (unsigned)TREE_INT_CST_LOW (gimple_call_arg (call, 5));
Ditto.
> +static void
> +oacc_loop_xform_head_tail (gcall *from, int level)
> +{
> + gimple_stmt_iterator gsi = gsi_for_stmt (from);
> + unsigned code = TREE_INT_CST_LOW (gimple_call_arg (from, 0));
> + tree replacement = build_int_cst (unsigned_type_node, level);
Too many spaces.
> + switch (gimple_call_internal_fn (call))
> + {
> + case IFN_UNIQUE:
> + {
> + unsigned c = TREE_INT_CST_LOW (gimple_call_arg (call, 0));
Shouldn't c be of type enum enum ifn_unique_kind ?
What about code?
> +
> + default:
> + break;
> + }
> + }
> +
> + break2:;
Can't you replace goto break2; with return; and
remove break2:; ?
> + if (TREE_INT_CST_LOW (gimple_call_arg (call, 0))
> + == IFN_GOACC_LOOP_BOUND)
> + goto break2;
> + }
> +
> + /* If we didn't see LOOP_BOUND, it should be in the single
> + successor block. */
> + basic_block bb = single_succ (gsi_bb (gsi));
> + gsi = gsi_start_bb (bb);
> + }
> +
> + break2:;
Similarly.
> + if (gimple_vdef (call))
> + replace_uses_by (gimple_vdef (call),
> + gimple_vuse (call));
Why the line break in between the arguments? The line wouldn't be really
long.
Otherwise LGTM.
Jakub
