OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity
0cc218d42c
gcc/gcc/omp-oacc-kernels-decompose.cc
Jakub Jelinek 99dee82307 Update copyright years.
2021-01-04 10:26:59 +01:00

1546 lines
54 KiB
C++
Raw Blame History

/* Decompose OpenACC 'kernels' constructs into parts, a sequence of compute
constructs
Copyright (C) 2020-2021 Free Software Foundation, Inc.
This file is part of GCC.
GCC is free software; you can redistribute it and/or modify it under
the terms of the GNU General Public License as published by the Free
Software Foundation; either version 3, or (at your option) any later
version.
GCC is distributed in the hope that it will be useful, but WITHOUT ANY
WARRANTY; without even the implied warranty of MERCHANTABILITY or
FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
for more details.
You should have received a copy of the GNU General Public License
along with GCC; see the file COPYING3. If not see
<http://www.gnu.org/licenses/>. */
#include "config.h"
#include "system.h"
#include "coretypes.h"
#include "backend.h"
#include "target.h"
#include "tree.h"
#include "langhooks.h"
#include "gimple.h"
#include "tree-pass.h"
#include "cgraph.h"
#include "fold-const.h"
#include "gimplify.h"
#include "gimple-iterator.h"
#include "gimple-walk.h"
#include "gomp-constants.h"
#include "omp-general.h"
#include "diagnostic-core.h"
/* This preprocessing pass is run immediately before lower_omp. It decomposes
OpenACC 'kernels' constructs into parts, a sequence of compute constructs.
The translation is as follows:
- The entire 'kernels' region is turned into a 'data' region with clauses
taken from the 'kernels' region. New 'create' clauses are added for all
variables declared at the top level in the kernels region.
- Any loop nests annotated with an OpenACC 'loop' directive are wrapped in
a new compute construct.
- 'loop' directives without an explicit 'independent' or 'seq' clause
get an 'auto' clause added; other clauses are preserved on the loop
or moved to the new surrounding compute construct, as applicable.
- Any sequences of other code (non-loops, non-OpenACC 'loop's) are wrapped
in new "gang-single" compute construct: 'worker'/'vector' parallelism is
preserved, but 'num_gangs (1)' is enforced.
- Both points above only apply at the topmost level in the region, that
is, the transformation does not introduce new compute constructs inside
nested statement bodies. In particular, this means that a
gang-parallelizable loop inside an 'if' statement is made "gang-single".
- In order to make the host wait only once for the whole region instead
of once per device kernel launch, the new compute constructs are
annotated 'async'. Unless the original 'kernels' construct already was
marked 'async', the entire region ends with a 'wait' directive. If the
original 'kernels' construct was marked 'async', the synthesized 'async'
clauses use the original 'kernels' construct's 'async' argument
(possibly implicit).
*/
/*TODO Things are conceptually wrong here: 'loop' clauses may be hidden behind
'device_type', so we have to defer a lot of processing until we're in the
offloading compilation. "Fortunately", GCC doesn't support the OpenACC
'device_type' clause yet, so we get away that. */
/* Helper function for decompose_kernels_region_body. If STMT contains a
"top-level" OMP_FOR statement, returns a pointer to that statement;
returns NULL otherwise.
A "top-level" OMP_FOR statement is one that is possibly accompanied by
small snippets of setup code. Specifically, this function accepts an
OMP_FOR possibly wrapped in a singleton bind and a singleton try
statement to allow for a local loop variable, but not an OMP_FOR
statement nested in any other constructs. Alternatively, it accepts a
non-singleton bind containing only assignments and then an OMP_FOR
statement at the very end. The former style can be generated by the C
frontend, the latter by the Fortran frontend. */
static gimple *
top_level_omp_for_in_stmt (gimple *stmt)
{
if (gimple_code (stmt) == GIMPLE_OMP_FOR)
return stmt;
if (gimple_code (stmt) == GIMPLE_BIND)
{
gimple_seq body = gimple_bind_body (as_a <gbind *> (stmt));
if (gimple_seq_singleton_p (body))
{
/* Accept an OMP_FOR statement, or a try statement containing only
a single OMP_FOR. */
gimple *maybe_for_or_try = gimple_seq_first_stmt (body);
if (gimple_code (maybe_for_or_try) == GIMPLE_OMP_FOR)
return maybe_for_or_try;
else if (gimple_code (maybe_for_or_try) == GIMPLE_TRY)
{
gimple_seq try_body = gimple_try_eval (maybe_for_or_try);
if (!gimple_seq_singleton_p (try_body))
return NULL;
gimple *maybe_omp_for_stmt = gimple_seq_first_stmt (try_body);
if (gimple_code (maybe_omp_for_stmt) == GIMPLE_OMP_FOR)
return maybe_omp_for_stmt;
}
}
else
{
gimple_stmt_iterator gsi;
/* Accept only a block of optional assignments followed by an
OMP_FOR at the end. No other kinds of statements allowed. */
for (gsi = gsi_start (body); !gsi_end_p (gsi); gsi_next (&gsi))
{
gimple *body_stmt = gsi_stmt (gsi);
if (gimple_code (body_stmt) == GIMPLE_ASSIGN)
continue;
else if (gimple_code (body_stmt) == GIMPLE_OMP_FOR
&& gsi_one_before_end_p (gsi))
return body_stmt;
else
return NULL;
}
}
}
return NULL;
}
/* Helper for adjust_region_code: evaluate the statement at GSI_P. */
static tree
adjust_region_code_walk_stmt_fn (gimple_stmt_iterator *gsi_p,
bool *handled_ops_p,
struct walk_stmt_info *wi)
{
int *region_code = (int *) wi->info;
gimple *stmt = gsi_stmt (*gsi_p);
switch (gimple_code (stmt))
{
case GIMPLE_OMP_FOR:
{
tree clauses = gimple_omp_for_clauses (stmt);
if (omp_find_clause (clauses, OMP_CLAUSE_INDEPENDENT))
{
/* Explicit 'independent' clause. */
/* Keep going; recurse into loop body. */
break;
}
else if (omp_find_clause (clauses, OMP_CLAUSE_SEQ))
{
/* Explicit 'seq' clause. */
/* We'll "parallelize" if at some level a loop construct has been
marked up by the user as unparallelizable ('seq' clause; we'll
respect that in the later processing). Given that the user has
explicitly marked it up, this loop construct cannot be
performance-critical, and in this case it's also fine to
"parallelize" instead of "gang-single", because any outer or
inner loops may still exploit the available parallelism. */
/* Keep going; recurse into loop body. */
break;
}
else
{
/* Explicit or implicit 'auto' clause. */
/* The user would like this loop analyzed ('auto' clause) and
typically parallelized, but we don't have available yet the
compiler logic to analyze this, so can't parallelize it here, so
we'd very likely be running into a performance problem if we
were to execute this unparallelized, thus forward the whole loop
nest to 'parloops'. */
*region_code = GF_OMP_TARGET_KIND_OACC_KERNELS;
/* Terminate: final decision for this region. */
*handled_ops_p = true;
return integer_zero_node;
}
gcc_unreachable ();
}
case GIMPLE_COND:
case GIMPLE_GOTO:
case GIMPLE_SWITCH:
case GIMPLE_ASM:
case GIMPLE_TRANSACTION:
case GIMPLE_RETURN:
/* Statement that might constitute some looping/control flow pattern. */
/* The user would like this code analyzed (implicit inside a 'kernels'
region) and typically parallelized, but we don't have available yet
the compiler logic to analyze this, so can't parallelize it here, so
we'd very likely be running into a performance problem if we were to
execute this unparallelized, thus forward the whole thing to
'parloops'. */
*region_code = GF_OMP_TARGET_KIND_OACC_KERNELS;
/* Terminate: final decision for this region. */
*handled_ops_p = true;
return integer_zero_node;
default:
/* Keep going. */
break;
}
return NULL;
}
/* Adjust the REGION_CODE for the region in GS. */
static void
adjust_region_code (gimple_seq gs, int *region_code)
{
struct walk_stmt_info wi;
memset (&wi, 0, sizeof (wi));
wi.info = region_code;
walk_gimple_seq (gs, adjust_region_code_walk_stmt_fn, NULL, &wi);
}
/* Helper function for make_loops_gang_single for walking the tree. If the
statement indicated by GSI_P is an OpenACC for loop with a gang clause,
issue a warning and remove the clause. */
static tree
visit_loops_in_gang_single_region (gimple_stmt_iterator *gsi_p,
bool *handled_ops_p,
struct walk_stmt_info *)
{
*handled_ops_p = false;
gimple *stmt = gsi_stmt (*gsi_p);
switch (gimple_code (stmt))
{
case GIMPLE_OMP_FOR:
/*TODO Given the current 'adjust_region_code' algorithm, this is
actually... */
gcc_unreachable ();
{
tree clauses = gimple_omp_for_clauses (stmt);
tree prev_clause = NULL;
for (tree clause = clauses; clause; clause = OMP_CLAUSE_CHAIN (clause))
{
if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_GANG)
{
/* It makes no sense to have a 'gang' clause in a "gang-single"
region, so warn and remove it. */
warning_at (gimple_location (stmt), 0,
"conditionally executed loop in %<kernels%> region"
" will be executed by a single gang;"
" ignoring %<gang%> clause");
if (prev_clause != NULL)
OMP_CLAUSE_CHAIN (prev_clause) = OMP_CLAUSE_CHAIN (clause);
else
clauses = OMP_CLAUSE_CHAIN (clause);
break;
}
prev_clause = clause;
}
gimple_omp_for_set_clauses (stmt, clauses);
}
/* No need to recurse into nested statements; no loop nested inside
this loop can be gang-partitioned. */
sorry ("%<gang%> loop in %<gang-single%> region");
*handled_ops_p = true;
break;
default:
break;
}
return NULL;
}
/* Visit all nested OpenACC loops in the sequence indicated by GS. This
statement is expected to be inside a gang-single region. Issue a warning
for any loops inside it that have gang clauses and remove the clauses. */
static void
make_loops_gang_single (gimple_seq gs)
{
struct walk_stmt_info wi;
memset (&wi, 0, sizeof (wi));
walk_gimple_seq (gs, visit_loops_in_gang_single_region, NULL, &wi);
}
/* Construct a "gang-single" compute construct at LOC containing the STMTS.
Annotate with CLAUSES, which must not contain a 'num_gangs' clause, and an
additional 'num_gangs (1)' clause to force "gang-single" execution. */
static gimple *
make_region_seq (location_t loc, gimple_seq stmts,
tree num_gangs_clause,
tree num_workers_clause,
tree vector_length_clause,
tree clauses)
{
/* This correctly unshares the entire clause chain rooted here. */
clauses = unshare_expr (clauses);
dump_user_location_t loc_stmts_first = gimple_seq_first (stmts);
/* Figure out the region code for this region. */
/* Optimistic default: assume "setup code", no looping; thus not
performance-critical. */
int region_code = GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE;
adjust_region_code (stmts, &region_code);
if (region_code == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE)
{
if (dump_enabled_p ())
/*TODO MSG_MISSED_OPTIMIZATION? */
dump_printf_loc (MSG_NOTE, loc_stmts_first,
"beginning %<gang-single%> part"
" in OpenACC %<kernels%> region\n");
/* Synthesize a 'num_gangs (1)' clause. */
tree gang_single_clause = build_omp_clause (loc, OMP_CLAUSE_NUM_GANGS);
OMP_CLAUSE_OPERAND (gang_single_clause, 0) = integer_one_node;
OMP_CLAUSE_CHAIN (gang_single_clause) = clauses;
clauses = gang_single_clause;
/* Remove and issue warnings about gang clauses on any OpenACC
loops nested inside this sequentially executed statement. */
make_loops_gang_single (stmts);
}
else if (region_code == GF_OMP_TARGET_KIND_OACC_KERNELS)
{
if (dump_enabled_p ())
dump_printf_loc (MSG_NOTE, loc_stmts_first,
"beginning %<parloops%> part"
" in OpenACC %<kernels%> region\n");
/* As we're transforming a 'GF_OMP_TARGET_KIND_OACC_KERNELS' into another
'GF_OMP_TARGET_KIND_OACC_KERNELS', this isn't doing any of the clauses
mangling that 'make_region_loop_nest' is doing. */
/* Re-assemble the clauses stripped off earlier. */
if (num_gangs_clause != NULL)
{
tree c = unshare_expr (num_gangs_clause);
OMP_CLAUSE_CHAIN (c) = clauses;
clauses = c;
}
if (num_workers_clause != NULL)
{
tree c = unshare_expr (num_workers_clause);
OMP_CLAUSE_CHAIN (c) = clauses;
clauses = c;
}
if (vector_length_clause != NULL)
{
tree c = unshare_expr (vector_length_clause);
OMP_CLAUSE_CHAIN (c) = clauses;
clauses = c;
}
}
else
gcc_unreachable ();
/* Build the gang-single region. */
gimple *single_region = gimple_build_omp_target (NULL, region_code, clauses);
gimple_set_location (single_region, loc);
gbind *single_body = gimple_build_bind (NULL, stmts, make_node (BLOCK));
gimple_omp_set_body (single_region, single_body);
return single_region;
}
/* Helper function for make_region_loop_nest. Adds a 'num_gangs'
('num_workers', 'vector_length') clause to the given CLAUSES, either the one
from the parent compute construct (PARENT_CLAUSE) or a new one based on the
loop's own LOOP_CLAUSE ('gang (num: N)' or similar for 'worker' or 'vector'
clauses) with the given CLAUSE_CODE. Does nothing if neither PARENT_CLAUSE
nor LOOP_CLAUSE exist. Returns the new clauses. */
static tree
add_parent_or_loop_num_clause (tree parent_clause, tree loop_clause,
omp_clause_code clause_code, tree clauses)
{
if (parent_clause != NULL)
{
tree num_clause = unshare_expr (parent_clause);
OMP_CLAUSE_CHAIN (num_clause) = clauses;
clauses = num_clause;
}
else if (loop_clause != NULL)
{
/* The kernels region does not have a 'num_gangs' clause, but the loop
itself had a 'gang (num: N)' clause. Honor it by adding a
'num_gangs (N)' clause on the compute construct. */
tree num = OMP_CLAUSE_OPERAND (loop_clause, 0);
tree new_num_clause
= build_omp_clause (OMP_CLAUSE_LOCATION (loop_clause), clause_code);
OMP_CLAUSE_OPERAND (new_num_clause, 0) = num;
OMP_CLAUSE_CHAIN (new_num_clause) = clauses;
clauses = new_num_clause;
}
return clauses;
}
/* Helper for make_region_loop_nest, looking for 'worker (num: N)' or 'vector
(length: N)' clauses in nested loops. Removes the argument, transferring it
to the enclosing compute construct (via WI->INFO). If arguments within the
same loop nest conflict, emits a warning.
This function also decides whether to add an 'auto' clause on each of these
nested loops. */
struct adjust_nested_loop_clauses_wi_info
{
tree *loop_gang_clause_ptr;
tree *loop_worker_clause_ptr;
tree *loop_vector_clause_ptr;
};
static tree
adjust_nested_loop_clauses (gimple_stmt_iterator *gsi_p, bool *,
struct walk_stmt_info *wi)
{
struct adjust_nested_loop_clauses_wi_info *wi_info
= (struct adjust_nested_loop_clauses_wi_info *) wi->info;
gimple *stmt = gsi_stmt (*gsi_p);
if (gimple_code (stmt) == GIMPLE_OMP_FOR)
{
bool add_auto_clause = true;
tree loop_clauses = gimple_omp_for_clauses (stmt);
tree loop_clause = loop_clauses;
for (; loop_clause; loop_clause = OMP_CLAUSE_CHAIN (loop_clause))
{
tree *outer_clause_ptr = NULL;
switch (OMP_CLAUSE_CODE (loop_clause))
{
case OMP_CLAUSE_GANG:
outer_clause_ptr = wi_info->loop_gang_clause_ptr;
break;
case OMP_CLAUSE_WORKER:
outer_clause_ptr = wi_info->loop_worker_clause_ptr;
break;
case OMP_CLAUSE_VECTOR:
outer_clause_ptr = wi_info->loop_vector_clause_ptr;
break;
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_AUTO:
add_auto_clause = false;
default:
break;
}
if (outer_clause_ptr != NULL)
{
if (OMP_CLAUSE_OPERAND (loop_clause, 0) != NULL
&& *outer_clause_ptr == NULL)
{
/* Transfer the clause to the enclosing compute construct and
remove the numerical argument from the 'loop'. */
*outer_clause_ptr = unshare_expr (loop_clause);
OMP_CLAUSE_OPERAND (loop_clause, 0) = NULL;
}
else if (OMP_CLAUSE_OPERAND (loop_clause, 0) != NULL &&
OMP_CLAUSE_OPERAND (*outer_clause_ptr, 0) != NULL)
{
/* See if both of these are the same constant. If they
aren't, emit a warning. */
tree old_op = OMP_CLAUSE_OPERAND (*outer_clause_ptr, 0);
tree new_op = OMP_CLAUSE_OPERAND (loop_clause, 0);
if (!(cst_and_fits_in_hwi (old_op) &&
cst_and_fits_in_hwi (new_op) &&
int_cst_value (old_op) == int_cst_value (new_op)))
{
const char *clause_name
= omp_clause_code_name[OMP_CLAUSE_CODE (loop_clause)];
error_at (gimple_location (stmt),
"cannot honor conflicting %qs clause",
clause_name);
inform (OMP_CLAUSE_LOCATION (*outer_clause_ptr),
"location of the previous clause"
" in the same loop nest");
}
OMP_CLAUSE_OPERAND (loop_clause, 0) = NULL;
}
}
}
if (add_auto_clause)
{
tree auto_clause
= build_omp_clause (gimple_location (stmt), OMP_CLAUSE_AUTO);
OMP_CLAUSE_CHAIN (auto_clause) = loop_clauses;
gimple_omp_for_set_clauses (stmt, auto_clause);
}
}
return NULL;
}
/* Helper for make_region_loop_nest. Transform OpenACC 'kernels'/'loop'
construct clauses into OpenACC 'parallel'/'loop' construct ones. */
static tree
transform_kernels_loop_clauses (gimple *omp_for,
tree num_gangs_clause,
tree num_workers_clause,
tree vector_length_clause,
tree clauses)
{
/* If this loop in a kernels region does not have an explicit 'seq',
'independent', or 'auto' clause, we must give it an explicit 'auto'
clause.
We also check for 'gang (num: N)' clauses. These must not appear in
kernels regions that have their own 'num_gangs' clause. Otherwise, they
must be converted and put on the region; similarly for 'worker' and
'vector' clauses. */
bool add_auto_clause = true;
tree loop_gang_clause = NULL, loop_worker_clause = NULL,
loop_vector_clause = NULL;
tree loop_clauses = gimple_omp_for_clauses (omp_for);
for (tree loop_clause = loop_clauses;
loop_clause;
loop_clause = OMP_CLAUSE_CHAIN (loop_clause))
{
bool found_num_clause = false;
tree *clause_ptr, clause_to_check;
switch (OMP_CLAUSE_CODE (loop_clause))
{
case OMP_CLAUSE_GANG:
found_num_clause = true;
clause_ptr = &loop_gang_clause;
clause_to_check = num_gangs_clause;
break;
case OMP_CLAUSE_WORKER:
found_num_clause = true;
clause_ptr = &loop_worker_clause;
clause_to_check = num_workers_clause;
break;
case OMP_CLAUSE_VECTOR:
found_num_clause = true;
clause_ptr = &loop_vector_clause;
clause_to_check = vector_length_clause;
break;
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_AUTO:
add_auto_clause = false;
default:
break;
}
if (found_num_clause && OMP_CLAUSE_OPERAND (loop_clause, 0) != NULL)
{
if (clause_to_check)
{
const char *clause_name
= omp_clause_code_name[OMP_CLAUSE_CODE (loop_clause)];
const char *parent_clause_name
= omp_clause_code_name[OMP_CLAUSE_CODE (clause_to_check)];
error_at (OMP_CLAUSE_LOCATION (loop_clause),
"argument not permitted on %qs clause"
" in OpenACC %<kernels%> region with a %qs clause",
clause_name, parent_clause_name);
inform (OMP_CLAUSE_LOCATION (clause_to_check),
"location of OpenACC %<kernels%>");
}
/* Copy the 'gang (N)'/'worker (N)'/'vector (N)' clause to the
enclosing compute construct. */
*clause_ptr = unshare_expr (loop_clause);
OMP_CLAUSE_CHAIN (*clause_ptr) = NULL;
/* Leave a 'gang'/'worker'/'vector' clause on the 'loop', but without
argument. */
OMP_CLAUSE_OPERAND (loop_clause, 0) = NULL;
}
}
if (add_auto_clause)
{
tree auto_clause = build_omp_clause (gimple_location (omp_for),
OMP_CLAUSE_AUTO);
OMP_CLAUSE_CHAIN (auto_clause) = loop_clauses;
loop_clauses = auto_clause;
}
gimple_omp_for_set_clauses (omp_for, loop_clauses);
/* We must also recurse into the loop; it might contain nested loops having
their own 'worker (num: W)' or 'vector (length: V)' clauses. Turn these
into 'worker'/'vector' clauses on the compute construct. */
struct walk_stmt_info wi;
memset (&wi, 0, sizeof (wi));
struct adjust_nested_loop_clauses_wi_info wi_info;
wi_info.loop_gang_clause_ptr = &loop_gang_clause;
wi_info.loop_worker_clause_ptr = &loop_worker_clause;
wi_info.loop_vector_clause_ptr = &loop_vector_clause;
wi.info = &wi_info;
gimple *body = gimple_omp_body (omp_for);
walk_gimple_seq (body, adjust_nested_loop_clauses, NULL, &wi);
/* Check if there were conflicting numbers of workers or vector length. */
if (loop_gang_clause != NULL &&
OMP_CLAUSE_OPERAND (loop_gang_clause, 0) == NULL)
loop_gang_clause = NULL;
if (loop_worker_clause != NULL &&
OMP_CLAUSE_OPERAND (loop_worker_clause, 0) == NULL)
loop_worker_clause = NULL;
if (loop_vector_clause != NULL &&
OMP_CLAUSE_OPERAND (loop_vector_clause, 0) == NULL)
vector_length_clause = NULL;
/* If the kernels region had 'num_gangs', 'num_worker', 'vector_length'
clauses, add these to this new compute construct. */
clauses
= add_parent_or_loop_num_clause (num_gangs_clause, loop_gang_clause,
OMP_CLAUSE_NUM_GANGS, clauses);
clauses
= add_parent_or_loop_num_clause (num_workers_clause, loop_worker_clause,
OMP_CLAUSE_NUM_WORKERS, clauses);
clauses
= add_parent_or_loop_num_clause (vector_length_clause, loop_vector_clause,
OMP_CLAUSE_VECTOR_LENGTH, clauses);
return clauses;
}
/* Construct a possibly gang-parallel compute construct containing the STMT,
which must be identical to, or a bind containing, the loop OMP_FOR.
The NUM_GANGS_CLAUSE, NUM_WORKERS_CLAUSE, and VECTOR_LENGTH_CLAUSE are
optional clauses from the original kernels region and must not be contained
in the other CLAUSES. The newly created compute construct is annotated with
the optional NUM_GANGS_CLAUSE as well as the other CLAUSES. If there is no
NUM_GANGS_CLAUSE but the loop has a 'gang (num: N)' clause, that is
converted to a 'num_gangs (N)' clause on the new compute construct, and
similarly for 'worker' and 'vector' clauses.
The outermost loop gets an 'auto' clause unless there already is an
'seq'/'independent'/'auto' clause. Nested loops inside OMP_FOR are treated
similarly by the adjust_nested_loop_clauses function. */
static gimple *
make_region_loop_nest (gimple *omp_for, gimple_seq stmts,
tree num_gangs_clause,
tree num_workers_clause,
tree vector_length_clause,
tree clauses)
{
/* This correctly unshares the entire clause chain rooted here. */
clauses = unshare_expr (clauses);
/* Figure out the region code for this region. */
/* Optimistic default: assume that the loop nest is parallelizable
(essentially, no GIMPLE_OMP_FOR with (explicit or implicit) 'auto' clause,
and no un-annotated loops). */
int region_code = GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED;
adjust_region_code (stmts, &region_code);
if (region_code == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED)
{
if (dump_enabled_p ())
/* This is not MSG_OPTIMIZED_LOCATIONS, as we're just doing what the
user asked us to. */
dump_printf_loc (MSG_NOTE, omp_for,
"parallelized loop nest"
" in OpenACC %<kernels%> region\n");
clauses = transform_kernels_loop_clauses (omp_for,
num_gangs_clause,
num_workers_clause,
vector_length_clause,
clauses);
}
else if (region_code == GF_OMP_TARGET_KIND_OACC_KERNELS)
{
if (dump_enabled_p ())
dump_printf_loc (MSG_NOTE, omp_for,
"forwarded loop nest"
" in OpenACC %<kernels%> region"
" to %<parloops%> for analysis\n");
/* We're transforming one 'GF_OMP_TARGET_KIND_OACC_KERNELS' into another
'GF_OMP_TARGET_KIND_OACC_KERNELS', so don't have to
'transform_kernels_loop_clauses'. */
/* Re-assemble the clauses stripped off earlier. */
clauses
= add_parent_or_loop_num_clause (num_gangs_clause, NULL,
OMP_CLAUSE_NUM_GANGS, clauses);
clauses
= add_parent_or_loop_num_clause (num_workers_clause, NULL,
OMP_CLAUSE_NUM_WORKERS, clauses);
clauses
= add_parent_or_loop_num_clause (vector_length_clause, NULL,
OMP_CLAUSE_VECTOR_LENGTH, clauses);
}
else
gcc_unreachable ();
gimple *parallel_body_bind
= gimple_build_bind (NULL, stmts, make_node (BLOCK));
gimple *parallel_region
= gimple_build_omp_target (parallel_body_bind, region_code, clauses);
gimple_set_location (parallel_region, gimple_location (omp_for));
return parallel_region;
}
/* Eliminate any binds directly inside BIND by adding their statements to
BIND (i.e., modifying it in place), excluding binds that hold only an
OMP_FOR loop and associated setup/cleanup code. Recurse into binds but
not other statements. Return a chain of the local variables of eliminated
binds, i.e., the local variables found in nested binds. If
INCLUDE_TOPLEVEL_VARS is true, this also includes the variables belonging
to BIND itself. */
static tree
flatten_binds (gbind *bind, bool include_toplevel_vars = false)
{
tree vars = NULL, last_var = NULL;
if (include_toplevel_vars)
{
vars = gimple_bind_vars (bind);
last_var = vars;
}
gimple_seq new_body = NULL;
gimple_seq body_sequence = gimple_bind_body (bind);
gimple_stmt_iterator gsi, gsi_n;
for (gsi = gsi_start (body_sequence); !gsi_end_p (gsi); gsi = gsi_n)
{
/* Advance the iterator here because otherwise it would be invalidated
by moving statements below. */
gsi_n = gsi;
gsi_next (&gsi_n);
gimple *stmt = gsi_stmt (gsi);
/* Flatten bind statements, except the ones that contain only an
OpenACC for loop. */
if (gimple_code (stmt) == GIMPLE_BIND
&& !top_level_omp_for_in_stmt (stmt))
{
gbind *inner_bind = as_a <gbind *> (stmt);
/* Flatten recursively, and collect all variables. */
tree inner_vars = flatten_binds (inner_bind, true);
gimple_seq inner_sequence = gimple_bind_body (inner_bind);
if (flag_checking)
{
for (gimple_stmt_iterator inner_gsi = gsi_start (inner_sequence);
!gsi_end_p (inner_gsi);
gsi_next (&inner_gsi))
{
gimple *inner_stmt = gsi_stmt (inner_gsi);
gcc_assert (gimple_code (inner_stmt) != GIMPLE_BIND
|| top_level_omp_for_in_stmt (inner_stmt));
}
}
gimple_seq_add_seq (&new_body, inner_sequence);
/* Find the last variable; we will append others to it. */
while (last_var != NULL && TREE_CHAIN (last_var) != NULL)
last_var = TREE_CHAIN (last_var);
if (last_var != NULL)
{
TREE_CHAIN (last_var) = inner_vars;
last_var = inner_vars;
}
else
{
vars = inner_vars;
last_var = vars;
}
}
else
gimple_seq_add_stmt (&new_body, stmt);
}
/* Put the possibly transformed body back into the bind. */
gimple_bind_set_body (bind, new_body);
return vars;
}
/* Helper function for places where we construct data regions. Wraps the BODY
inside a try-finally construct at LOC that calls __builtin_GOACC_data_end
in its cleanup block. Returns this try statement. */
static gimple *
make_data_region_try_statement (location_t loc, gimple *body)
{
tree data_end_fn = builtin_decl_explicit (BUILT_IN_GOACC_DATA_END);
gimple *call = gimple_build_call (data_end_fn, 0);
gimple_seq cleanup = NULL;
gimple_seq_add_stmt (&cleanup, call);
gimple *try_stmt = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
gimple_set_location (body, loc);
return try_stmt;
}
/* If INNER_BIND_VARS holds variables, build an OpenACC data region with
location LOC containing BODY and having 'create (var)' clauses for each
variable. If INNER_CLEANUP is present, add a try-finally statement with
this cleanup code in the finally block. Return the new data region, or
the original BODY if no data region was needed. */
static gimple *
maybe_build_inner_data_region (location_t loc, gimple *body,
tree inner_bind_vars, gimple *inner_cleanup)
{
/* Is this an instantiation of a template? (In this case, we don't care what
the generic decl is - just whether the function decl has one.) */
bool generic_inst_p
= (lang_hooks.decls.get_generic_function_decl (current_function_decl)
!= NULL);
/* Build data 'create (var)' clauses for these local variables.
Below we will add these to a data region enclosing the entire body
of the decomposed kernels region. */
tree prev_mapped_var = NULL, next = NULL, artificial_vars = NULL,
inner_data_clauses = NULL;
for (tree v = inner_bind_vars; v; v = next)
{
next = TREE_CHAIN (v);
if (DECL_ARTIFICIAL (v)
|| TREE_CODE (v) == CONST_DECL
|| generic_inst_p)
{
/* If this is an artificial temporary, it need not be mapped. We
move its declaration into the bind inside the data region.
Also avoid mapping variables if we are inside a template
instantiation; the code does not contain all the copies to
temporaries that would make this legal. */
TREE_CHAIN (v) = artificial_vars;
artificial_vars = v;
if (prev_mapped_var != NULL)
TREE_CHAIN (prev_mapped_var) = next;
else
inner_bind_vars = next;
}
else
{
/* Otherwise, build the map clause. */
tree new_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (new_clause, GOMP_MAP_ALLOC);
OMP_CLAUSE_DECL (new_clause) = v;
OMP_CLAUSE_SIZE (new_clause) = DECL_SIZE_UNIT (v);
OMP_CLAUSE_CHAIN (new_clause) = inner_data_clauses;
inner_data_clauses = new_clause;
prev_mapped_var = v;
}
}
if (artificial_vars)
body = gimple_build_bind (artificial_vars, body, make_node (BLOCK));
/* If we determined above that there are variables that need to be created
on the device, construct a data region for them and wrap the body
inside that. */
if (inner_data_clauses != NULL)
{
gcc_assert (inner_bind_vars != NULL);
gimple *inner_data_region
= gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DATA_KERNELS,
inner_data_clauses);
gimple_set_location (inner_data_region, loc);
/* Make sure __builtin_GOACC_data_end is called at the end. */
gimple *try_stmt = make_data_region_try_statement (loc, body);
gimple_omp_set_body (inner_data_region, try_stmt);
gimple *bind_body;
if (inner_cleanup != NULL)
/* Clobber all the inner variables that need to be clobbered. */
bind_body = gimple_build_try (inner_data_region, inner_cleanup,
GIMPLE_TRY_FINALLY);
else
bind_body = inner_data_region;
body = gimple_build_bind (inner_bind_vars, bind_body, make_node (BLOCK));
}
return body;
}
/* Helper function of decompose_kernels_region_body. The statements in
REGION_BODY are expected to be decomposed parts; add an 'async' clause to
each. Also add a 'wait' directive at the end of the sequence. */
static void
add_async_clauses_and_wait (location_t loc, gimple_seq *region_body)
{
tree default_async_queue
= build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL);
for (gimple_stmt_iterator gsi = gsi_start (*region_body);
!gsi_end_p (gsi);
gsi_next (&gsi))
{
gimple *stmt = gsi_stmt (gsi);
tree target_clauses = gimple_omp_target_clauses (stmt);
tree new_async_clause = build_omp_clause (loc, OMP_CLAUSE_ASYNC);
OMP_CLAUSE_OPERAND (new_async_clause, 0) = default_async_queue;
OMP_CLAUSE_CHAIN (new_async_clause) = target_clauses;
target_clauses = new_async_clause;
gimple_omp_target_set_clauses (as_a <gomp_target *> (stmt),
target_clauses);
}
/* A '#pragma acc wait' is just a call 'GOACC_wait (acc_async_sync, 0)'. */
tree wait_fn = builtin_decl_explicit (BUILT_IN_GOACC_WAIT);
tree sync_arg = build_int_cst (integer_type_node, GOMP_ASYNC_SYNC);
gimple *wait_call = gimple_build_call (wait_fn, 2,
sync_arg, integer_zero_node);
gimple_set_location (wait_call, loc);
gimple_seq_add_stmt (region_body, wait_call);
}
/* Auxiliary analysis of the body of a kernels region, to determine for each
OpenACC loop whether it is control-dependent (i.e., not necessarily
executed every time the kernels region is entered) or not.
We say that a loop is control-dependent if there is some cond, switch, or
goto statement that jumps over it, forwards or backwards. For example,
if the loop is controlled by an if statement, then a jump to the true
block, the false block, or from one of those blocks to the control flow
join point will necessarily jump over the loop.
This analysis implements an ad-hoc union-find data structure classifying
statements into "control-flow regions" as follows: Most statements are in
the same region as their predecessor, except that each OpenACC loop is in
a region of its own, and each OpenACC loop's successor starts a new
region. We then unite the regions of any statements linked by jumps,
placing any cond, switch, or goto statement in the same region as its
target label(s).
In the end, control dependence of OpenACC loops can be determined by
comparing their immediate predecessor and successor statements' regions.
A jump crosses the loop if and only if the predecessor and successor are
in the same region. (If there is no predecessor or successor, the loop
is executed unconditionally.)
The methods in this class identify statements by their index in the
kernels region's body. */
class control_flow_regions
{
public:
/* Initialize an instance and pre-compute the control-flow region
information for the statement sequence SEQ. */
control_flow_regions (gimple_seq seq);
/* Return true if the statement with the given index IDX in the analyzed
statement sequence is an unconditionally executed OpenACC loop. */
bool is_unconditional_oacc_for_loop (size_t idx);
private:
/* Find the region representative for the statement identified by index
STMT_IDX. */
size_t find_rep (size_t stmt_idx);
/* Union the regions containing the statements represented by
representatives A and B. */
void union_reps (size_t a, size_t b);
/* Helper for the constructor. Performs the actual computation of the
control-flow regions in the statement sequence SEQ. */
void compute_regions (gimple_seq seq);
/* The mapping from statement indices to region representatives. */
vec <size_t> representatives;
/* A cache mapping statement indices to a flag indicating whether the
statement is a top level OpenACC for loop. */
vec <bool> omp_for_loops;
};
control_flow_regions::control_flow_regions (gimple_seq seq)
{
representatives.create (1);
omp_for_loops.create (1);
compute_regions (seq);
}
bool
control_flow_regions::is_unconditional_oacc_for_loop (size_t idx)
{
if (idx == 0 || idx == representatives.length () - 1)
/* The first or last statement in the kernels region. This means that
there is no room before or after it for a jump or a label. Thus
there cannot be a jump across it, so it is unconditional. */
return true;
/* Otherwise, the loop is unconditional if the statements before and after
it are in different control flow regions. Scan forward and backward,
skipping over neighboring OpenACC for loops, to find these preceding
statements. */
size_t prev_index = idx - 1;
while (prev_index > 0 && omp_for_loops [prev_index] == true)
prev_index--;
/* If all preceding statements are also OpenACC loops, all of these are
unconditional. */
if (prev_index == 0)
return true;
size_t succ_index = idx + 1;
while (succ_index < omp_for_loops.length ()
&& omp_for_loops [succ_index] == true)
succ_index++;
/* If all following statements are also OpenACC loops, all of these are
unconditional. */
if (succ_index == omp_for_loops.length ())
return true;
return (find_rep (prev_index) != find_rep (succ_index));
}
size_t
control_flow_regions::find_rep (size_t stmt_idx)
{
size_t rep = stmt_idx, aux = stmt_idx;
/* Find the root representative of this statement. */
while (representatives[rep] != rep)
rep = representatives[rep];
/* Compress the path from the original statement to the representative. */
while (representatives[aux] != rep)
{
size_t tmp = representatives[aux];
representatives[aux] = rep;
aux = tmp;
}
return rep;
}
void
control_flow_regions::union_reps (size_t a, size_t b)
{
a = find_rep (a);
b = find_rep (b);
representatives[b] = a;
}
void
control_flow_regions::compute_regions (gimple_seq seq)
{
hash_map <gimple *, size_t> control_flow_reps;
hash_map <tree, size_t> label_reps;
size_t current_region = 0, idx = 0;
/* In a first pass, assign an initial region to each statement. Except in
the case of OpenACC loops, each statement simply gets the same region
representative as its predecessor. */
for (gimple_stmt_iterator gsi = gsi_start (seq);
!gsi_end_p (gsi);
gsi_next (&gsi))
{
gimple *stmt = gsi_stmt (gsi);
gimple *omp_for = top_level_omp_for_in_stmt (stmt);
omp_for_loops.safe_push (omp_for != NULL);
if (omp_for != NULL)
{
/* Assign a new region to this loop and to its successor. */
current_region = idx;
representatives.safe_push (current_region);
current_region++;
}
else
{
representatives.safe_push (current_region);
/* Remember any jumps and labels for the second pass below. */
if (gimple_code (stmt) == GIMPLE_COND
|| gimple_code (stmt) == GIMPLE_SWITCH
|| gimple_code (stmt) == GIMPLE_GOTO)
control_flow_reps.put (stmt, current_region);
else if (gimple_code (stmt) == GIMPLE_LABEL)
label_reps.put (gimple_label_label (as_a <glabel *> (stmt)),
current_region);
}
idx++;
}
gcc_assert (representatives.length () == omp_for_loops.length ());
/* Revisit all the control flow statements and union the region of each
cond, switch, or goto statement with the target labels' regions. */
for (hash_map <gimple *, size_t>::iterator it = control_flow_reps.begin ();
it != control_flow_reps.end ();
++it)
{
gimple *stmt = (*it).first;
size_t stmt_rep = (*it).second;
switch (gimple_code (stmt))
{
tree label;
unsigned int n;
case GIMPLE_COND:
label = gimple_cond_true_label (as_a <gcond *> (stmt));
union_reps (stmt_rep, *label_reps.get (label));
label = gimple_cond_false_label (as_a <gcond *> (stmt));
union_reps (stmt_rep, *label_reps.get (label));
break;
case GIMPLE_SWITCH:
n = gimple_switch_num_labels (as_a <gswitch *> (stmt));
for (unsigned int i = 0; i < n; i++)
{
tree switch_case
= gimple_switch_label (as_a <gswitch *> (stmt), i);
label = CASE_LABEL (switch_case);
union_reps (stmt_rep, *label_reps.get (label));
}
break;
case GIMPLE_GOTO:
label = gimple_goto_dest (stmt);
union_reps (stmt_rep, *label_reps.get (label));
break;
default:
gcc_unreachable ();
}
}
}
/* Decompose the body of the KERNELS_REGION, which was originally annotated
with the KERNELS_CLAUSES, into a series of compute constructs. */
static gimple *
decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
{
location_t loc = gimple_location (kernels_region);
/* The kernels clauses will be propagated to the child clauses unmodified,
except that the 'num_gangs', 'num_workers', and 'vector_length' clauses
will only be added to loop regions. The other regions are "gang-single"
and get an explicit 'num_gangs (1)' clause. So separate out the
'num_gangs', 'num_workers', and 'vector_length' clauses here.
Also check for the presence of an 'async' clause but do not remove it from
the 'kernels' clauses. */
tree num_gangs_clause = NULL, num_workers_clause = NULL,
vector_length_clause = NULL;
tree async_clause = NULL;
tree prev_clause = NULL, next_clause = NULL;
tree parallel_clauses = kernels_clauses;
for (tree c = parallel_clauses; c; c = next_clause)
{
/* Preserve this here, as we might NULL it later. */
next_clause = OMP_CLAUSE_CHAIN (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_GANGS
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_WORKERS
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR_LENGTH)
{
/* Cut this clause out of the chain. */
if (prev_clause != NULL)
OMP_CLAUSE_CHAIN (prev_clause) = OMP_CLAUSE_CHAIN (c);
else
kernels_clauses = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = NULL;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_NUM_GANGS:
num_gangs_clause = c;
break;
case OMP_CLAUSE_NUM_WORKERS:
num_workers_clause = c;
break;
case OMP_CLAUSE_VECTOR_LENGTH:
vector_length_clause = c;
break;
default:
gcc_unreachable ();
}
}
else
prev_clause = c;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
async_clause = c;
}
gimple *kernels_body = gimple_omp_body (kernels_region);
gbind *kernels_bind = as_a <gbind *> (kernels_body);
/* The body of the region may contain other nested binds declaring inner
local variables. Collapse all these binds into one to ensure that we
have a single sequence of statements to iterate over; also, collect all
inner variables. */
tree inner_bind_vars = flatten_binds (kernels_bind);
gimple_seq body_sequence = gimple_bind_body (kernels_bind);
/* All these inner variables will get allocated on the device (below, by
calling maybe_build_inner_data_region). Here we create 'present'
clauses for them and add these clauses to the list of clauses to be
attached to each inner compute construct. */
tree present_clauses = kernels_clauses;
for (tree var = inner_bind_vars; var; var = TREE_CHAIN (var))
{
if (!DECL_ARTIFICIAL (var) && TREE_CODE (var) != CONST_DECL)
{
tree present_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (present_clause, GOMP_MAP_FORCE_PRESENT);
OMP_CLAUSE_DECL (present_clause) = var;
OMP_CLAUSE_SIZE (present_clause) = DECL_SIZE_UNIT (var);
OMP_CLAUSE_CHAIN (present_clause) = present_clauses;
present_clauses = present_clause;
}
}
kernels_clauses = present_clauses;
/* In addition to nested binds, the "real" body of the region may be
nested inside a try-finally block. Find its cleanup block, which
contains code to clobber the local variables that must be clobbered. */
gimple *inner_cleanup = NULL;
if (body_sequence != NULL && gimple_code (body_sequence) == GIMPLE_TRY)
{
if (gimple_seq_singleton_p (body_sequence))
{
/* The try statement is the only thing inside the bind. */
inner_cleanup = gimple_try_cleanup (body_sequence);
body_sequence = gimple_try_eval (body_sequence);
}
else
{
/* The bind's body starts with a try statement, but it is followed
by other things. */
gimple_stmt_iterator gsi = gsi_start (body_sequence);
gimple *try_stmt = gsi_stmt (gsi);
inner_cleanup = gimple_try_cleanup (try_stmt);
gimple *try_body = gimple_try_eval (try_stmt);
gsi_remove (&gsi, false);
/* Now gsi indicates the sequence of statements after the try
statement in the bind. Append the statement in the try body and
the trailing statements from gsi. */
gsi_insert_seq_before (&gsi, try_body, GSI_CONTINUE_LINKING);
body_sequence = gsi_stmt (gsi);
}
}
/* This sequence will collect all the top-level statements in the body of
the data region we are about to construct. */
gimple_seq region_body = NULL;
/* This sequence will collect consecutive statements to be put into a
gang-single region. */
gimple_seq gang_single_seq = NULL;
/* Flag recording whether the gang_single_seq only contains copies to
local variables. These may be loop setup code that should not be
separated from the loop. */
bool only_simple_assignments = true;
/* Precompute the control flow region information to determine whether an
OpenACC loop is executed conditionally or unconditionally. */
control_flow_regions cf_regions (body_sequence);
/* Iterate over the statements in the kernels region's body. */
size_t idx = 0;
gimple_stmt_iterator gsi, gsi_n;
for (gsi = gsi_start (body_sequence); !gsi_end_p (gsi); gsi = gsi_n, idx++)
{
/* Advance the iterator here because otherwise it would be invalidated
by moving statements below. */
gsi_n = gsi;
gsi_next (&gsi_n);
gimple *stmt = gsi_stmt (gsi);
gimple *omp_for = top_level_omp_for_in_stmt (stmt);
bool is_unconditional_oacc_for_loop = false;
if (omp_for != NULL)
is_unconditional_oacc_for_loop
= cf_regions.is_unconditional_oacc_for_loop (idx);
if (omp_for != NULL
&& is_unconditional_oacc_for_loop)
{
/* This is an OMP for statement, put it into a separate region.
But first, construct a gang-single region containing any
complex sequential statements we may have seen. */
if (gang_single_seq != NULL && !only_simple_assignments)
{
gimple *single_region
= make_region_seq (loc, gang_single_seq,
num_gangs_clause,
num_workers_clause,
vector_length_clause,
kernels_clauses);
gimple_seq_add_stmt (&region_body, single_region);
}
else if (gang_single_seq != NULL && only_simple_assignments)
{
/* There is a sequence of sequential statements preceding this
loop, but they are all simple assignments. This is
probably setup code for the loop; in particular, Fortran DO
loops are preceded by code to copy the loop limit variable
to a temporary. Group this code together with the loop
itself. */
gimple_seq_add_stmt (&gang_single_seq, stmt);
stmt = gimple_build_bind (NULL, gang_single_seq,
make_node (BLOCK));
}
gang_single_seq = NULL;
only_simple_assignments = true;
gimple_seq parallel_seq = NULL;
gimple_seq_add_stmt (&parallel_seq, stmt);
gimple *parallel_region
= make_region_loop_nest (omp_for, parallel_seq,
num_gangs_clause,
num_workers_clause,
vector_length_clause,
kernels_clauses);
gimple_seq_add_stmt (&region_body, parallel_region);
}
else
{
if (omp_for != NULL)
{
gcc_checking_assert (!is_unconditional_oacc_for_loop);
if (dump_enabled_p ())
dump_printf_loc (MSG_MISSED_OPTIMIZATION, omp_for,
"unparallelized loop nest"
" in OpenACC %<kernels%> region:"
" it's executed conditionally\n");
}
/* This is not an unconditional OMP for statement, so it will be
put into a gang-single region. */
gimple_seq_add_stmt (&gang_single_seq, stmt);
/* Is this a simple assignment? We call it simple if it is an
assignment to an artificial local variable. This captures
Fortran loop setup code computing loop bounds and offsets. */
bool is_simple_assignment
= (gimple_code (stmt) == GIMPLE_ASSIGN
&& TREE_CODE (gimple_assign_lhs (stmt)) == VAR_DECL
&& DECL_ARTIFICIAL (gimple_assign_lhs (stmt)));
if (!is_simple_assignment)
only_simple_assignments = false;
}
}
/* If we did not emit a new region, and are not going to emit one now
(that is, the original region was empty), prepare to emit a dummy so as
to preserve the original construct, which other processing (at least
test cases) depend on. */
if (region_body == NULL && gang_single_seq == NULL)
{
gimple *stmt = gimple_build_nop ();
gimple_set_location (stmt, loc);
gimple_seq_add_stmt (&gang_single_seq, stmt);
}
/* Gather up any remaining gang-single statements. */
if (gang_single_seq != NULL)
{
gimple *single_region
= make_region_seq (loc, gang_single_seq,
num_gangs_clause,
num_workers_clause,
vector_length_clause,
kernels_clauses);
gimple_seq_add_stmt (&region_body, single_region);
}
/* We want to launch these kernels asynchronously. If the original
kernels region had an async clause, this is done automatically because
that async clause was copied to the individual regions we created.
Otherwise, add an async clause to each newly created region, as well as
a wait directive at the end. */
if (async_clause == NULL)
add_async_clauses_and_wait (loc, &region_body);
tree kernels_locals = gimple_bind_vars (as_a <gbind *> (kernels_body));
gimple *body = gimple_build_bind (kernels_locals, region_body,
make_node (BLOCK));
/* If we found variables declared in nested scopes, build a data region to
map them to the device. */
body = maybe_build_inner_data_region (loc, body, inner_bind_vars,
inner_cleanup);
return body;
}
/* Decompose one OpenACC 'kernels' construct into an OpenACC 'data' construct
containing the original OpenACC 'kernels' construct's region cut up into a
sequence of compute constructs. */
static gimple *
omp_oacc_kernels_decompose_1 (gimple *kernels_stmt)
{
gcc_checking_assert (gimple_omp_target_kind (kernels_stmt)
== GF_OMP_TARGET_KIND_OACC_KERNELS);
location_t loc = gimple_location (kernels_stmt);
/* Collect the data clauses of the OpenACC 'kernels' directive and create a
new OpenACC 'data' construct with those clauses. */
tree kernels_clauses = gimple_omp_target_clauses (kernels_stmt);
tree data_clauses = NULL;
for (tree c = kernels_clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
/* Certain clauses are copied to the enclosing OpenACC 'data'. Other
clauses remain on the OpenACC 'kernels'. */
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
{
tree decl = OMP_CLAUSE_DECL (c);
HOST_WIDE_INT map_kind = OMP_CLAUSE_MAP_KIND (c);
switch (map_kind)
{
default:
if (map_kind == GOMP_MAP_ALLOC
&& integer_zerop (OMP_CLAUSE_SIZE (c)))
/* ??? This is an alloc clause for mapping a pointer whose
target is already mapped. We leave these on the inner
compute constructs because moving them to the outer data
region causes runtime errors. */
break;
/* For non-artificial variables, and for non-declaration
expressions like A[0:n], copy the clause to the data
region. */
if ((DECL_P (decl) && !DECL_ARTIFICIAL (decl))
|| !DECL_P (decl))
{
tree new_clause = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (new_clause, map_kind);
/* This must be unshared here to avoid "incorrect sharing
of tree nodes" errors from verify_gimple. */
OMP_CLAUSE_DECL (new_clause) = unshare_expr (decl);
OMP_CLAUSE_SIZE (new_clause) = OMP_CLAUSE_SIZE (c);
OMP_CLAUSE_CHAIN (new_clause) = data_clauses;
data_clauses = new_clause;
/* Now that this data is mapped, turn the data clause on the
inner OpenACC 'kernels' into a 'present' clause. */
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT);
}
break;
case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
/* ??? Copying these map kinds leads to internal compiler
errors in later passes. */
break;
}
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF)
{
/* If there is an 'if' clause, it must be duplicated to the
enclosing data region. Temporarily remove the if clause's
chain to avoid copying it. */
tree saved_chain = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = NULL;
tree new_if_clause = unshare_expr (c);
OMP_CLAUSE_CHAIN (c) = saved_chain;
OMP_CLAUSE_CHAIN (new_if_clause) = data_clauses;
data_clauses = new_if_clause;
}
}
/* Restore the original order of the clauses. */
data_clauses = nreverse (data_clauses);
gimple *data_region
= gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DATA_KERNELS,
data_clauses);
gimple_set_location (data_region, loc);
/* Transform the body of the kernels region into a sequence of compute
constructs. */
gimple *body = decompose_kernels_region_body (kernels_stmt,
kernels_clauses);
/* Put the transformed pieces together. The entire body of the region is
wrapped in a try-finally statement that calls __builtin_GOACC_data_end
for cleanup. */
gimple *try_stmt = make_data_region_try_statement (loc, body);
gimple_omp_set_body (data_region, try_stmt);
return data_region;
}
/* Decompose OpenACC 'kernels' constructs in the current function. */
static tree
omp_oacc_kernels_decompose_callback_stmt (gimple_stmt_iterator *gsi_p,
bool *handled_ops_p,
struct walk_stmt_info *)
{
gimple *stmt = gsi_stmt (*gsi_p);
if ((gimple_code (stmt) == GIMPLE_OMP_TARGET)
&& gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS)
{
gimple *stmt_new = omp_oacc_kernels_decompose_1 (stmt);
gsi_replace (gsi_p, stmt_new, false);
*handled_ops_p = true;
}
else
*handled_ops_p = false;
return NULL;
}
static unsigned int
omp_oacc_kernels_decompose (void)
{
gimple_seq body = gimple_body (current_function_decl);
struct walk_stmt_info wi;
memset (&wi, 0, sizeof (wi));
walk_gimple_seq_mod (&body, omp_oacc_kernels_decompose_callback_stmt, NULL,
&wi);
gimple_set_body (current_function_decl, body);
return 0;
}
namespace {
const pass_data pass_data_omp_oacc_kernels_decompose =
{
GIMPLE_PASS, /* type */
"omp_oacc_kernels_decompose", /* name */
OPTGROUP_OMP, /* optinfo_flags */
TV_NONE, /* tv_id */
PROP_gimple_any, /* properties_required */
0, /* properties_provided */
0, /* properties_destroyed */
0, /* todo_flags_start */
0, /* todo_flags_finish */
};
class pass_omp_oacc_kernels_decompose : public gimple_opt_pass
{
public:
pass_omp_oacc_kernels_decompose (gcc::context *ctxt)
: gimple_opt_pass (pass_data_omp_oacc_kernels_decompose, ctxt)
{}
/* opt_pass methods: */
virtual bool gate (function *)
{
return (flag_openacc
&& flag_openacc_kernels == OPENACC_KERNELS_DECOMPOSE);
}
virtual unsigned int execute (function *)
{
return omp_oacc_kernels_decompose ();
}
}; // class pass_omp_oacc_kernels_decompose
} // anon namespace
gimple_opt_pass *
make_pass_omp_oacc_kernels_decompose (gcc::context *ctxt)
{
return new pass_omp_oacc_kernels_decompose (ctxt);
}
Reference in New Issue View Git Blame Copy Permalink