This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4] Accelerator constructs omp lowering and expansion
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Richard Henderson <rth at redhat dot com>, "Michael V. Zolotukhin" <michael dot v dot zolotukhin at gmail dot com>
- Cc: gcc-patches at gcc dot gnu dot org
- Date: Wed, 4 Sep 2013 20:54:47 +0200
- Subject: [gomp4] Accelerator constructs omp lowering and expansion
- Authentication-results: sourceware.org; auth=none
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
Hi!
This patch implements #pragma omp {target{, data, update},teams} lowering
and expansion, and adds stub calls into libgomp, so that (for now
unconditionally) we can at least always fall back to host execution.
I still need to add testcases and test reference based array section
handling and look at mapping of VLAs.
Any issues you can spot in the patch? If not, I'll commit it tomorrow, so
that I can continue poking in these areas without too much pending stuff.
Known issues that don't affect host fallback, but will be a problem
for offloading (Richard or others, any thoughts about that?):
1) right now we pass a host fn address and a function name string
to GOMP_target. For host execution we of course only need the fn
address, for offloading I guess libgomp will need to dlinfo
that fn address to find out what shared library (or binary) contains
it and somehow look for the special data section in it.
But, right now the *._omp_fn.N functions are always static, so there
is nothing to look up by name for Intel MIC, and if we wanted to make
the symbol global, we'd need to give it a shlib resp. binary unique
name; but if the containing function is not global, how can we do that?
Append get_file_function_name to the name?
2) much bigger problem seems to be global #pragma omp declare target
variables. Those are supposed to be mapped from the beginning,
if they are just copied into the target LTO subset streaming, they will
be emitted normally as data variables. But, unfortunately the
runtime must be aware of those mappings, because you can do stuff like:
#pragma omp declare target
int v = 1;
#pragma omp end declare target
void bar (int *p)
{
#pragma omp target map(to:p[:1])
*p++;
}
void foo () { bar (&v); }
where the runtime should assign target's p copy the value of target's
v variable. Or even for say #pragma omp target update.
So, on the host side, we need to prepare the triplets of host var
address, var size and during linking somehow supply it info on how
to create the target address, and have some function (one per
shared library resp. binary) that would locate the data section
within the shared library/binary for the requested accelerator
and with another argument call some libgomp function to initialize
the device data environment for the given shared library/binary
and accelerator. Perhaps we could pass the address of such (.hidden)
function as yet another argument to GOMP_target call and
--as-needed link it from some *.a library?
2013-09-04 Jakub Jelinek <jakub@redhat.com>
* tree-cfg.c (make_edges): For GIMPLE_OMP_TARGET
with GF_OMP_TARGET_KIND_UPDATE, don't look for
GIMPLE_OMP_RETURN and immediately restore previous
region.
* langhooks.c (lhd_omp_mappable_type): New function.
* omp-low.c (scan_sharing_clauses): Ignore OMP_CLAUSE_SHARED
in GIMPLE_OMP_TEAMS constructs. Handle OMP_CLAUSE_NUM_TEAMS,
OMP_CLAUSE_THREAD_LIMIT, OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP,
OMP_CLAUSE_TO and OMP_CLAUSE_FROM.
(create_omp_child_function): If current function has
"omp declare target" attribute or if current region
is OMP_TARGET or lexically nested in it, add that
attribute to the omp child function.
(scan_omp_target, scan_omp_teams): New functions.
(check_omp_nesting_restrictions): Fix a typo in TEAMS nesting
check.
(scan_omp_1_stmt): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
(lower_rec_input_clauses): Ignore OMP_CLAUSE_SHARED
in GIMPLE_OMP_TEAMS constructs.
(expand_omp_synch): Handle GIMPLE_OMP_TEAMS.
(expand_omp_target): New function.
(expand_omp): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
(build_omp_regions_1): For GIMPLE_OMP_TARGET with
GF_OMP_TARGET_KIND_UPDATE, don't look for GIMPLE_OMP_RETURN and
immediately restore previous region.
(lower_omp_single): Emit a CLOBBER stmt after GIMPLE_OMP_RETURN.
(lower_omp_taskreg): Likewise.
(lower_omp_target, lower_omp_teams): New functions.
(lower_omp_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
* tree.h (enum omp_clause_map_kind): Improve description of
OMP_CLAUSE_MAP_POINTER.
(OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION): Define.
* gimple.def (GIMPLE_OMP_TARGET): Use GSS_OMP_PARALLEL instead
of GSS_OMP_SINGLE.
* langhooks.h (struct lang_hooks_for_types): Add
omp_mappable_type hook.
* langhooks-def.h (lhd_omp_mappable_type): New prototype.
(LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
(LANG_HOOKS_FOR_TYPES_INITIALIZER): Use it.
* gimple.h (gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
gimple_omp_target_set_clauses): Use gimple_omp_parallel
instead of gimple_omp_single.
(gimple_omp_target_child_fn, gimple_omp_target_child_fn_ptr,
gimple_omp_target_set_child_fn, gimple_omp_target_data_arg,
gimple_omp_target_data_arg_ptr, gimple_omp_target_set_data_arg): New
inlines.
* omp-builtins.def (BUILT_IN_GOMP_TARGET, BUILT_IN_GOMP_TARGET_DATA,
BUILT_IN_GOMP_TARGET_END_DATA, BUILT_IN_GOMP_TARGET_UPDATE,
BUILT_IN_GOMP_TEAMS): New builtins.
* gimple-pretty-print.c (dump_gimple_omp_target): Print child_fn
for #pragma omp target.
* gimplify.c (omp_notice_variable): Diagnose if implicitly mapped
decl doesn't have mappable type.
(gimplify_scan_omp_clauses): For OMP_CLAUSE_MAP don't clear
notice_outer.
(gimplify_adjust_omp_clauses_1): Add OMP_CLAUSE_MAP even if the same
decl is already referenced in target data construct surrounding it.
(gimplify_adjust_omp_clauses): Likewise. Handle
OMP_CLAUSE_THREAD_LIMIT.
(gimplify_omp_workshare): Fix up gimplification of target or target
data construct body. For target data add GOMP_target_end_data
call in a try/finally cleanup.
* builtin-types.def (BT_FN_VOID_UINT_UINT,
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
c/
* c-typeck.c (handle_omp_array_sections_1): Remove pointer_based_p
argument. Do c_save_expr on low bound unconditionally.
(handle_omp_array_sections): Adjust caller. Call c_fully_fold when
needed. Add OMP_CLAUSE_MAP_POINTER map clause even for array based
array sections.
(c_finish_omp_clauses): Don't complain about non-mappable types
for OMP_CLAUSE_MAP_POINTER. Check for duplicates even for map
clauses. Handle OMP_CLAUSE_NUM_TEAMS.
* c-parser.c (c_parser_omp_target_data, c_parser_omp_target): Call
keep_next_level ().
cp/
* parser.c (cp_parser_omp_target_data, cp_parser_omp_target): Call
keep_next_level (true).
* cp-objcp-common.h (LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
* semantics.c (handle_omp_array_sections_1): Remove pointer_based_p
argument. Do cp_save_expr on low bound unconditionally.
(handle_omp_array_sections): Adjust caller. Add OMP_CLAUSE_MAP_POINTER
map clause even for array based array sections.
(finish_omp_clauses): Don't complain about non-mappable types
for OMP_CLAUSE_MAP_POINTER. Check for duplicates even for map
clauses.
fortran/
* types.def (BT_FN_VOID_UINT_UINT, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
testsuite/
* c-c++-common/gomp/map-1.c (foo): Add dg-error for implicitly
mapped non-mappable var.
libgomp/
* libgomp.map (GOMP_4.0): Add GOMP_target, GOMP_target_data,
GOMP_target_end_data, GOMP_target_update and GOMP_teams.
* Makefile.am (libgomp_la_SOURCES): Add target.c.
* Makefile.in: Regenerated.
* target.c: New file.
* libgomp_g.h (GOMP_target, GOMP_target_data,
GOMP_target_end_data, GOMP_target_update, GOMP_teams): New prototypes.
* testsuite/libgomp.c++/for-11.C (main): Uncomment
#pragma omp target teams directive.
* testsuite/libgomp.c/for-3.c (main): Likewise.
* testsuite/libgomp.c++/target-1.C: New test.
* testsuite/libgomp.c/target-1.c: New test.
--- gcc/tree-cfg.c.jj 2013-08-27 20:53:36.000000000 +0200
+++ gcc/tree-cfg.c 2013-09-02 13:55:36.591737508 +0200
@@ -610,7 +610,6 @@ make_edges (void)
case GIMPLE_OMP_TASK:
case GIMPLE_OMP_FOR:
case GIMPLE_OMP_SINGLE:
- case GIMPLE_OMP_TARGET:
case GIMPLE_OMP_TEAMS:
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
@@ -620,6 +619,13 @@ make_edges (void)
fallthru = true;
break;
+ case GIMPLE_OMP_TARGET:
+ cur_region = new_omp_region (bb, code, cur_region);
+ fallthru = true;
+ if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE)
+ cur_region = cur_region->outer;
+ break;
+
case GIMPLE_OMP_SECTIONS:
cur_region = new_omp_region (bb, code, cur_region);
fallthru = true;
--- gcc/langhooks.c.jj 2013-08-27 20:53:03.000000000 +0200
+++ gcc/langhooks.c 2013-09-04 19:30:55.551423852 +0200
@@ -523,6 +523,15 @@ lhd_omp_firstprivatize_type_sizes (struc
{
}
+/* Return true if TYPE is an OpenMP mappable type. By default return true
+ if type is complete. */
+
+bool
+lhd_omp_mappable_type (tree type)
+{
+ return COMPLETE_TYPE_P (type);
+}
+
/* Common function for add_builtin_function and
add_builtin_function_ext_scope. */
static tree
--- gcc/omp-low.c.jj 2013-08-27 22:44:31.000000000 +0200
+++ gcc/omp-low.c 2013-09-04 19:58:30.320019227 +0200
@@ -1453,6 +1453,9 @@ scan_sharing_clauses (tree clauses, omp_
break;
case OMP_CLAUSE_SHARED:
+ /* Ignore shared directives in teams construct. */
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+ break;
gcc_assert (is_taskreg_ctx (ctx));
decl = OMP_CLAUSE_DECL (c);
gcc_assert (!COMPLETE_TYPE_P (TREE_TYPE (decl))
@@ -1533,6 +1536,9 @@ scan_sharing_clauses (tree clauses, omp_
case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
+ case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_THREAD_LIMIT:
+ case OMP_CLAUSE_DEVICE:
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEPEND:
@@ -1540,6 +1546,66 @@ scan_sharing_clauses (tree clauses, omp_
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ case OMP_CLAUSE_MAP:
+ if (ctx->outer)
+ scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
+ decl = OMP_CLAUSE_DECL (c);
+ /* Global variables with "omp declare target" attribute
+ don't need to be copied, the receiver side will use them
+ directly. */
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && DECL_P (decl)
+ && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl)))
+ break;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ {
+ /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
+ #pragma omp target data, there is nothing to map for
+ those. */
+ if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
+ && !POINTER_TYPE_P (TREE_TYPE (decl)))
+ break;
+ }
+ if (DECL_P (decl))
+ {
+ install_var_field (decl, true, 3, ctx);
+ if (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_REGION)
+ install_var_local (decl, ctx);
+ }
+ else
+ {
+ tree base = get_base_address (decl);
+ tree nc = OMP_CLAUSE_CHAIN (c);
+ if (DECL_P (base)
+ && nc != NULL_TREE
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_DECL (nc) == base
+ && OMP_CLAUSE_MAP_KIND (nc) == OMP_CLAUSE_MAP_POINTER
+ && integer_zerop (OMP_CLAUSE_SIZE (nc)))
+ {
+ OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) = 1;
+ OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (nc) = 1;
+ }
+ else
+ {
+ gcc_assert (!splay_tree_lookup (ctx->field_map,
+ (splay_tree_key) decl));
+ tree field
+ = build_decl (OMP_CLAUSE_LOCATION (c),
+ FIELD_DECL, NULL_TREE, ptr_type_node);
+ insert_field_into_struct (ctx->record_type, field);
+ splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
+ (splay_tree_value) field);
+ }
+ }
+ break;
+
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_COLLAPSE:
@@ -1590,16 +1656,46 @@ scan_sharing_clauses (tree clauses, omp_
break;
case OMP_CLAUSE_SHARED:
+ /* Ignore shared directives in teams construct. */
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+ break;
decl = OMP_CLAUSE_DECL (c);
if (! is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
fixup_remapped_decl (decl, ctx, false);
break;
+ case OMP_CLAUSE_MAP:
+ if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+ break;
+ decl = OMP_CLAUSE_DECL (c);
+ if (DECL_P (decl)
+ && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl)))
+ break;
+ if (DECL_P (decl))
+ {
+ if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
+ && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
+ {
+ tree new_decl = lookup_decl (decl, ctx);
+ TREE_TYPE (new_decl)
+ = remap_type (TREE_TYPE (decl), &ctx->cb);
+ }
+ else
+ fixup_remapped_decl (decl, ctx, false);
+ }
+ break;
+
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_DEFAULT:
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
+ case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_THREAD_LIMIT:
+ case OMP_CLAUSE_DEVICE:
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_NOWAIT:
@@ -1613,6 +1709,8 @@ scan_sharing_clauses (tree clauses, omp_
case OMP_CLAUSE_ALIGNED:
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE__LOOPTEMP_:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
break;
default:
@@ -1677,6 +1775,26 @@ create_omp_child_function (omp_context *
DECL_EXTERNAL (decl) = 0;
DECL_CONTEXT (decl) = NULL_TREE;
DECL_INITIAL (decl) = make_node (BLOCK);
+ bool target_p = false;
+ if (lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (current_function_decl)))
+ target_p = true;
+ else
+ {
+ omp_context *octx;
+ for (octx = ctx; octx; octx = octx->outer)
+ if (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET
+ && gimple_omp_target_kind (octx->stmt)
+ == GF_OMP_TARGET_KIND_REGION)
+ {
+ target_p = true;
+ break;
+ }
+ }
+ if (target_p)
+ DECL_ATTRIBUTES (decl)
+ = tree_cons (get_identifier ("omp declare target"),
+ NULL_TREE, DECL_ATTRIBUTES (decl));
t = build_decl (DECL_SOURCE_LOCATION (decl),
RESULT_DECL, NULL_TREE, void_type_node);
@@ -1975,6 +2093,53 @@ scan_omp_single (gimple stmt, omp_contex
layout_type (ctx->record_type);
}
+/* Scan an OpenMP target{, data, update} directive. */
+
+static void
+scan_omp_target (gimple stmt, omp_context *outer_ctx)
+{
+ omp_context *ctx;
+ tree name;
+ int kind = gimple_omp_target_kind (stmt);
+
+ ctx = new_omp_context (stmt, outer_ctx);
+ ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
+ ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
+ ctx->record_type = lang_hooks.types.make_type (RECORD_TYPE);
+ name = create_tmp_var_name (".omp_data_t");
+ name = build_decl (gimple_location (stmt),
+ TYPE_DECL, name, ctx->record_type);
+ DECL_ARTIFICIAL (name) = 1;
+ DECL_NAMELESS (name) = 1;
+ TYPE_NAME (ctx->record_type) = name;
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ create_omp_child_function (ctx, false);
+ gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
+ }
+
+ scan_sharing_clauses (gimple_omp_target_clauses (stmt), ctx);
+ scan_omp (gimple_omp_body_ptr (stmt), ctx);
+
+ if (TYPE_FIELDS (ctx->record_type) == NULL)
+ ctx->record_type = ctx->receiver_decl = NULL;
+ else
+ {
+ layout_type (ctx->record_type);
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ fixup_child_record_type (ctx);
+ }
+}
+
+/* Scan an OpenMP teams directive. */
+
+static void
+scan_omp_teams (gimple stmt, omp_context *outer_ctx)
+{
+ omp_context *ctx = new_omp_context (stmt, outer_ctx);
+ scan_sharing_clauses (gimple_omp_teams_clauses (stmt), ctx);
+ scan_omp (gimple_omp_body_ptr (stmt), ctx);
+}
/* Check OpenMP nesting restrictions. */
static bool
@@ -1992,7 +2157,7 @@ check_omp_nesting_restrictions (gimple s
else if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
{
if ((gimple_code (stmt) != GIMPLE_OMP_FOR
- || (gimple_omp_for_kind (ctx->stmt)
+ || (gimple_omp_for_kind (stmt)
!= GF_OMP_FOR_KIND_DISTRIBUTE))
&& gimple_code (stmt) != GIMPLE_OMP_PARALLEL)
{
@@ -2345,6 +2510,14 @@ scan_omp_1_stmt (gimple_stmt_iterator *g
scan_omp (gimple_omp_body_ptr (stmt), ctx);
break;
+ case GIMPLE_OMP_TARGET:
+ scan_omp_target (stmt, ctx);
+ break;
+
+ case GIMPLE_OMP_TEAMS:
+ scan_omp_teams (stmt, ctx);
+ break;
+
case GIMPLE_BIND:
{
tree var;
@@ -2731,6 +2904,9 @@ lower_rec_input_clauses (tree clauses, g
continue;
break;
case OMP_CLAUSE_SHARED:
+ /* Ignore shared directives in teams construct. */
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+ continue;
if (maybe_lookup_decl (OMP_CLAUSE_DECL (c), ctx) == NULL)
{
gcc_assert (is_global_var (OMP_CLAUSE_DECL (c)));
@@ -2889,6 +3065,9 @@ lower_rec_input_clauses (tree clauses, g
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_SHARED:
+ /* Ignore shared directives in teams construct. */
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
+ continue;
/* Shared global vars are just accessed directly. */
if (is_global_var (new_var))
break;
@@ -6728,7 +6907,8 @@ expand_omp_synch (struct omp_region *reg
gcc_assert (gimple_code (gsi_stmt (si)) == GIMPLE_OMP_SINGLE
|| gimple_code (gsi_stmt (si)) == GIMPLE_OMP_MASTER
|| gimple_code (gsi_stmt (si)) == GIMPLE_OMP_ORDERED
- || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_CRITICAL);
+ || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_CRITICAL
+ || gimple_code (gsi_stmt (si)) == GIMPLE_OMP_TEAMS);
gsi_remove (&si, true);
single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
@@ -7317,6 +7497,318 @@ expand_omp_atomic (struct omp_region *re
}
+/* Expand the OpenMP target{, data, update} directive starting at REGION. */
+
+static void
+expand_omp_target (struct omp_region *region)
+{
+ basic_block entry_bb, exit_bb, new_bb;
+ struct function *child_cfun = NULL;
+ tree child_fn = NULL_TREE, block, t;
+ gimple_stmt_iterator gsi;
+ gimple entry_stmt, stmt;
+ edge e;
+
+ entry_stmt = last_stmt (region->entry);
+ new_bb = region->entry;
+ int kind = gimple_omp_target_kind (entry_stmt);
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ child_fn = gimple_omp_target_child_fn (entry_stmt);
+ child_cfun = DECL_STRUCT_FUNCTION (child_fn);
+ }
+
+ entry_bb = region->entry;
+ exit_bb = region->exit;
+
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ unsigned srcidx, dstidx, num;
+
+ /* If the target region needs data sent from the parent
+ function, then the very first statement (except possible
+ tree profile counter updates) of the parallel body
+ is a copy assignment .OMP_DATA_I = &.OMP_DATA_O. Since
+ &.OMP_DATA_O is passed as an argument to the child function,
+ we need to replace it with the argument as seen by the child
+ function.
+
+ In most cases, this will end up being the identity assignment
+ .OMP_DATA_I = .OMP_DATA_I. However, if the parallel body had
+ a function call that has been inlined, the original PARM_DECL
+ .OMP_DATA_I may have been converted into a different local
+ variable. In which case, we need to keep the assignment. */
+ if (gimple_omp_target_data_arg (entry_stmt))
+ {
+ basic_block entry_succ_bb = single_succ (entry_bb);
+ gimple_stmt_iterator gsi;
+ tree arg;
+ gimple tgtcopy_stmt = NULL;
+ tree sender
+ = TREE_VEC_ELT (gimple_omp_target_data_arg (entry_stmt), 0);
+
+ for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
+ {
+ gcc_assert (!gsi_end_p (gsi));
+ stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) != GIMPLE_ASSIGN)
+ continue;
+
+ if (gimple_num_ops (stmt) == 2)
+ {
+ tree arg = gimple_assign_rhs1 (stmt);
+
+ /* We're ignoring the subcode because we're
+ effectively doing a STRIP_NOPS. */
+
+ if (TREE_CODE (arg) == ADDR_EXPR
+ && TREE_OPERAND (arg, 0) == sender)
+ {
+ tgtcopy_stmt = stmt;
+ break;
+ }
+ }
+ }
+
+ gcc_assert (tgtcopy_stmt != NULL);
+ arg = DECL_ARGUMENTS (child_fn);
+
+ gcc_assert (gimple_assign_lhs (tgtcopy_stmt) == arg);
+ gsi_remove (&gsi, true);
+ }
+
+ /* Declare local variables needed in CHILD_CFUN. */
+ block = DECL_INITIAL (child_fn);
+ BLOCK_VARS (block) = vec2chain (child_cfun->local_decls);
+ /* The gimplifier could record temporaries in target block
+ rather than in containing function's local_decls chain,
+ which would mean cgraph missed finalizing them. Do it now. */
+ for (t = BLOCK_VARS (block); t; t = DECL_CHAIN (t))
+ if (TREE_CODE (t) == VAR_DECL
+ && TREE_STATIC (t)
+ && !DECL_EXTERNAL (t))
+ varpool_finalize_decl (t);
+ DECL_SAVED_TREE (child_fn) = NULL;
+ /* We'll create a CFG for child_fn, so no gimple body is needed. */
+ gimple_set_body (child_fn, NULL);
+ TREE_USED (block) = 1;
+
+ /* Reset DECL_CONTEXT on function arguments. */
+ for (t = DECL_ARGUMENTS (child_fn); t; t = DECL_CHAIN (t))
+ DECL_CONTEXT (t) = child_fn;
+
+ /* Split ENTRY_BB at GIMPLE_OMP_TARGET,
+ so that it can be moved to the child function. */
+ gsi = gsi_last_bb (entry_bb);
+ stmt = gsi_stmt (gsi);
+ gcc_assert (stmt && gimple_code (stmt) == GIMPLE_OMP_TARGET
+ && gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_REGION);
+ gsi_remove (&gsi, true);
+ e = split_block (entry_bb, stmt);
+ entry_bb = e->dest;
+ single_succ_edge (entry_bb)->flags = EDGE_FALLTHRU;
+
+ /* Convert GIMPLE_OMP_RETURN into a RETURN_EXPR. */
+ if (exit_bb)
+ {
+ gsi = gsi_last_bb (exit_bb);
+ gcc_assert (!gsi_end_p (gsi)
+ && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+ stmt = gimple_build_return (NULL);
+ gsi_insert_after (&gsi, stmt, GSI_SAME_STMT);
+ gsi_remove (&gsi, true);
+ }
+
+ /* Move the target region into CHILD_CFUN. */
+
+ block = gimple_block (entry_stmt);
+
+ new_bb = move_sese_region_to_fn (child_cfun, entry_bb, exit_bb, block);
+ if (exit_bb)
+ single_succ_edge (new_bb)->flags = EDGE_FALLTHRU;
+ /* When the OMP expansion process cannot guarantee an up-to-date
+ loop tree arrange for the child function to fixup loops. */
+ if (loops_state_satisfies_p (LOOPS_NEED_FIXUP))
+ child_cfun->x_current_loops->state |= LOOPS_NEED_FIXUP;
+
+ /* Remove non-local VAR_DECLs from child_cfun->local_decls list. */
+ num = vec_safe_length (child_cfun->local_decls);
+ for (srcidx = 0, dstidx = 0; srcidx < num; srcidx++)
+ {
+ t = (*child_cfun->local_decls)[srcidx];
+ if (DECL_CONTEXT (t) == cfun->decl)
+ continue;
+ if (srcidx != dstidx)
+ (*child_cfun->local_decls)[dstidx] = t;
+ dstidx++;
+ }
+ if (dstidx != num)
+ vec_safe_truncate (child_cfun->local_decls, dstidx);
+
+ /* Inform the callgraph about the new function. */
+ DECL_STRUCT_FUNCTION (child_fn)->curr_properties = cfun->curr_properties;
+ cgraph_add_new_function (child_fn, true);
+
+ /* Fix the callgraph edges for child_cfun. Those for cfun will be
+ fixed in a following pass. */
+ push_cfun (child_cfun);
+ rebuild_cgraph_edges ();
+
+ /* Some EH regions might become dead, see PR34608. If
+ pass_cleanup_cfg isn't the first pass to happen with the
+ new child, these dead EH edges might cause problems.
+ Clean them up now. */
+ if (flag_exceptions)
+ {
+ basic_block bb;
+ bool changed = false;
+
+ FOR_EACH_BB (bb)
+ changed |= gimple_purge_dead_eh_edges (bb);
+ if (changed)
+ cleanup_tree_cfg ();
+ }
+ pop_cfun ();
+ }
+
+ /* Emit a library call to launch the target region, or do data
+ transfers. */
+ tree t1, t2, t3, t4, device, cond, c, clauses;
+ enum built_in_function start_ix;
+ location_t clause_loc;
+
+ clauses = gimple_omp_target_clauses (entry_stmt);
+
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ start_ix = BUILT_IN_GOMP_TARGET;
+ else if (kind == GF_OMP_TARGET_KIND_DATA)
+ start_ix = BUILT_IN_GOMP_TARGET_DATA;
+ else
+ start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+
+ /* By default, the value of DEVICE is -1 (let runtime library choose)
+ and there is no conditional. */
+ cond = NULL_TREE;
+ device = build_int_cst (integer_type_node, -1);
+
+ c = find_omp_clause (clauses, OMP_CLAUSE_IF);
+ if (c)
+ cond = OMP_CLAUSE_IF_EXPR (c);
+
+ c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+ if (c)
+ {
+ device = OMP_CLAUSE_DEVICE_ID (c);
+ clause_loc = OMP_CLAUSE_LOCATION (c);
+ }
+ else
+ clause_loc = gimple_location (entry_stmt);
+
+ /* Ensure 'device' is of the correct type. */
+ device = fold_convert_loc (clause_loc, integer_type_node, device);
+
+ /* If we found the clause 'if (cond)', build
+ (cond ? device : -2). */
+ if (cond)
+ {
+ cond = gimple_boolify (cond);
+
+ basic_block cond_bb, then_bb, else_bb;
+ edge e;
+ tree tmp_var;
+
+ tmp_var = create_tmp_var (TREE_TYPE (device), NULL);
+ e = split_block (entry_bb, NULL);
+ cond_bb = e->src;
+ entry_bb = e->dest;
+ remove_edge (e);
+
+ then_bb = create_empty_bb (cond_bb);
+ else_bb = create_empty_bb (then_bb);
+ set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+ set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb);
+
+ stmt = gimple_build_cond_empty (cond);
+ gsi = gsi_start_bb (cond_bb);
+ gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+ gsi = gsi_start_bb (then_bb);
+ stmt = gimple_build_assign (tmp_var, device);
+ gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+ gsi = gsi_start_bb (else_bb);
+ stmt = gimple_build_assign (tmp_var,
+ build_int_cst (integer_type_node, -2));
+ gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+ make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+ make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
+ if (current_loops)
+ {
+ add_bb_to_loop (then_bb, cond_bb->loop_father);
+ add_bb_to_loop (else_bb, cond_bb->loop_father);
+ }
+ make_edge (then_bb, entry_bb, EDGE_FALLTHRU);
+ make_edge (else_bb, entry_bb, EDGE_FALLTHRU);
+
+ device = tmp_var;
+ }
+
+ gsi = gsi_last_bb (new_bb);
+ t = gimple_omp_target_data_arg (entry_stmt);
+ if (t == NULL)
+ {
+ t1 = size_zero_node;
+ t2 = build_zero_cst (ptr_type_node);
+ t3 = t2;
+ t4 = t2;
+ }
+ else
+ {
+ t1 = TYPE_MAX_VALUE (TYPE_DOMAIN (TREE_TYPE (TREE_VEC_ELT (t, 1))));
+ t1 = size_binop (PLUS_EXPR, t1, size_int (1));
+ t2 = build_fold_addr_expr (TREE_VEC_ELT (t, 0));
+ t3 = build_fold_addr_expr (TREE_VEC_ELT (t, 1));
+ t4 = build_fold_addr_expr (TREE_VEC_ELT (t, 2));
+ }
+
+ gimple g;
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ tree fnaddr = build_fold_addr_expr (child_fn);
+ unsigned fnnamelen = IDENTIFIER_LENGTH (DECL_NAME (child_fn));
+ tree fnname = build_string (fnnamelen,
+ IDENTIFIER_POINTER (DECL_NAME (child_fn)));
+ TREE_TYPE (fnname) = build_array_type_nelts (char_type_node,
+ fnnamelen);
+ TREE_READONLY (fnname) = 1;
+ TREE_STATIC (fnname) = 1;
+ fnname = build_fold_addr_expr (fnname);
+ g = gimple_build_call (builtin_decl_explicit (start_ix), 7,
+ device, fnaddr, fnname, t1, t2, t3, t4);
+ }
+ else
+ g = gimple_build_call (builtin_decl_explicit (start_ix), 5,
+ device, t1, t2, t3, t4);
+ gimple_set_location (g, gimple_location (entry_stmt));
+ gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+ if (kind != GF_OMP_TARGET_KIND_REGION)
+ {
+ g = gsi_stmt (gsi);
+ gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
+ gsi_remove (&gsi, true);
+ }
+ if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+ {
+ gsi = gsi_last_bb (region->exit);
+ g = gsi_stmt (gsi);
+ gcc_assert (g && gimple_code (g) == GIMPLE_OMP_RETURN);
+ gsi_remove (&gsi, true);
+ }
+}
+
+
/* Expand the parallel region tree rooted at REGION. Expansion
proceeds in depth-first order. Innermost regions are expanded
first. This way, parallel regions that require a new function to
@@ -7374,6 +7866,7 @@ expand_omp (struct omp_region *region)
case GIMPLE_OMP_MASTER:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
+ case GIMPLE_OMP_TEAMS:
expand_omp_synch (region);
break;
@@ -7381,6 +7874,10 @@ expand_omp (struct omp_region *region)
expand_omp_atomic (region);
break;
+ case GIMPLE_OMP_TARGET:
+ expand_omp_target (region);
+ break;
+
default:
gcc_unreachable ();
}
@@ -7445,6 +7942,9 @@ build_omp_regions_1 (basic_block bb, str
GIMPLE_OMP_SECTIONS, and we do nothing for it. */
;
}
+ else if (code == GIMPLE_OMP_TARGET
+ && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE)
+ new_omp_region (bb, code, parent);
else
{
/* Otherwise, this directive becomes the parent for a new
@@ -7825,7 +8325,7 @@ lower_omp_single (gimple_stmt_iterator *
{
tree block;
gimple t, bind, single_stmt = gsi_stmt (*gsi_p);
- gimple_seq bind_body, dlist;
+ gimple_seq bind_body, bind_body_tail = NULL, dlist;
struct gimplify_ctx gctx;
push_gimplify_context (&gctx);
@@ -7855,8 +8355,17 @@ lower_omp_single (gimple_stmt_iterator *
t = gimple_build_omp_return
(!!find_omp_clause (gimple_omp_single_clauses (single_stmt),
OMP_CLAUSE_NOWAIT));
- gimple_seq_add_stmt (&bind_body, t);
- maybe_add_implicit_barrier_cancel (ctx, &bind_body);
+ gimple_seq_add_stmt (&bind_body_tail, t);
+ maybe_add_implicit_barrier_cancel (ctx, &bind_body_tail);
+ 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;
+ gsi_insert_after (&gsi, gimple_build_assign (ctx->sender_decl,
+ clobber), GSI_SAME_STMT);
+ }
+ gimple_seq_add_seq (&bind_body, bind_body_tail);
gimple_bind_set_body (bind, bind_body);
pop_gimplify_context (bind);
@@ -8605,6 +9114,14 @@ lower_omp_taskreg (gimple_stmt_iterator
lower_send_clauses (clauses, &ilist, &olist, ctx);
lower_send_shared_vars (&ilist, &olist, ctx);
+ if (ctx->record_type)
+ {
+ tree clobber = build_constructor (TREE_TYPE (ctx->sender_decl), NULL);
+ TREE_THIS_VOLATILE (clobber) = 1;
+ gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
+ clobber));
+ }
+
/* Once all the expansions are done, sequence all the different
fragments inside gimple_omp_body. */
@@ -8637,6 +9154,326 @@ lower_omp_taskreg (gimple_stmt_iterator
pop_gimplify_context (NULL);
}
+/* Lower the OpenMP target directive in the current statement
+ in GSI_P. CTX holds context information for the directive. */
+
+static void
+lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+ tree clauses;
+ tree child_fn, t, c;
+ gimple stmt = gsi_stmt (*gsi_p);
+ gimple tgt_bind = NULL, bind;
+ gimple_seq tgt_body = NULL, olist, ilist, new_body;
+ struct gimplify_ctx gctx;
+ location_t loc = gimple_location (stmt);
+ int kind = gimple_omp_target_kind (stmt);
+ unsigned int map_cnt = 0;
+
+ clauses = gimple_omp_target_clauses (stmt);
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
+ tgt_body = gimple_bind_body (tgt_bind);
+ }
+ else if (kind == GF_OMP_TARGET_KIND_DATA)
+ tgt_body = gimple_omp_body (stmt);
+ child_fn = ctx->cb.dst_fn;
+
+ push_gimplify_context (&gctx);
+
+ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ tree var, x;
+
+ default:
+ break;
+ case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ var = OMP_CLAUSE_DECL (c);
+ if (!DECL_P (var))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ map_cnt++;
+ continue;
+ }
+ if (!lookup_sfield (var, ctx))
+ continue;
+
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ x = build_receiver_ref (var, true, ctx);
+ tree new_var = lookup_decl (var, ctx);
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ map_cnt++;
+ }
+
+ if (kind != GF_OMP_TARGET_KIND_UPDATE)
+ lower_omp (&tgt_body, ctx);
+
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ /* Declare all the variables created by mapping and the variables
+ declared in the scope of the target body. */
+ record_vars_into (ctx->block_vars, child_fn);
+ record_vars_into (gimple_bind_vars (tgt_bind), child_fn);
+ }
+
+ olist = NULL;
+ ilist = NULL;
+ if (ctx->record_type)
+ {
+ ctx->sender_decl
+ = create_tmp_var (ctx->record_type, ".omp_data_arr");
+ DECL_NAMELESS (ctx->sender_decl) = 1;
+ TREE_ADDRESSABLE (ctx->sender_decl) = 1;
+ t = make_tree_vec (3);
+ TREE_VEC_ELT (t, 0) = ctx->sender_decl;
+ TREE_VEC_ELT (t, 1)
+ = create_tmp_var (build_array_type_nelts (size_type_node, map_cnt),
+ ".omp_data_sizes");
+ DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
+ TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
+ TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
+ TREE_VEC_ELT (t, 2)
+ = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
+ map_cnt),
+ ".omp_data_kinds");
+ DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
+ TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
+ TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
+ gimple_omp_target_set_data_arg (stmt, t);
+
+ vec<constructor_elt, va_gc> *vsize;
+ vec<constructor_elt, va_gc> *vkind;
+ vec_alloc (vsize, map_cnt);
+ vec_alloc (vkind, map_cnt);
+ unsigned int map_idx = 0;
+
+ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ tree ovar, nc;
+
+ default:
+ break;
+ case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ nc = c;
+ ovar = OMP_CLAUSE_DECL (c);
+ if (!DECL_P (ovar))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ {
+ gcc_checking_assert (OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (c))
+ == get_base_address (ovar));
+ nc = OMP_CLAUSE_CHAIN (c);
+ ovar = OMP_CLAUSE_DECL (nc);
+ }
+ else
+ {
+ tree x = build_sender_ref (ovar, ctx);
+ tree v
+ = build_fold_addr_expr_with_type (ovar, ptr_type_node);
+ gimplify_assign (x, v, &ilist);
+ nc = NULL_TREE;
+ }
+ }
+ else if (!lookup_sfield (ovar, ctx))
+ continue;
+
+ if (nc)
+ {
+ tree var = lookup_decl_in_outer_ctx (ovar, ctx);
+ tree x = build_sender_ref (ovar, ctx);
+ if (is_gimple_reg (var))
+ {
+ gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
+ tree avar = create_tmp_var (TREE_TYPE (var), NULL);
+ mark_addressable (avar);
+ if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
+ && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+ gimplify_assign (avar, var, &ilist);
+ avar = build_fold_addr_expr (avar);
+ gimplify_assign (x, avar, &ilist);
+ if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+ && !TYPE_READONLY (TREE_TYPE (var)))
+ {
+ x = build_sender_ref (ovar, ctx);
+ x = build_simple_mem_ref (x);
+ gimplify_assign (var, x, &olist);
+ }
+ }
+ else
+ {
+ var = build_fold_addr_expr (var);
+ gimplify_assign (x, var, &ilist);
+ }
+ }
+ tree s = OMP_CLAUSE_SIZE (c);
+ if (s == NULL_TREE)
+ s = TYPE_SIZE (TREE_TYPE (ovar));
+ s = fold_convert (size_type_node, s);
+ tree purpose = size_int (map_idx++);
+ CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
+ if (TREE_CODE (s) != INTEGER_CST)
+ TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
+
+ unsigned char tkind = 0;
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_MAP:
+ tkind = OMP_CLAUSE_MAP_KIND (c);
+ break;
+ case OMP_CLAUSE_TO:
+ tkind = OMP_CLAUSE_MAP_TO;
+ break;
+ case OMP_CLAUSE_FROM:
+ tkind = OMP_CLAUSE_MAP_FROM;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+ build_int_cst (unsigned_char_type_node,
+ tkind));
+ if (nc && nc != c)
+ c = nc;
+ }
+
+ gcc_assert (map_idx == map_cnt);
+
+ DECL_INITIAL (TREE_VEC_ELT (t, 1))
+ = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
+ DECL_INITIAL (TREE_VEC_ELT (t, 2))
+ = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind);
+ if (!TREE_STATIC (TREE_VEC_ELT (t, 1)))
+ {
+ gimple_seq initlist = NULL;
+ force_gimple_operand (build1 (DECL_EXPR, void_type_node,
+ TREE_VEC_ELT (t, 1)),
+ &initlist, true, NULL_TREE);
+ gimple_seq_add_seq (&ilist, initlist);
+ }
+
+ tree clobber = build_constructor (ctx->record_type, NULL);
+ TREE_THIS_VOLATILE (clobber) = 1;
+ gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
+ clobber));
+ }
+
+ /* Once all the expansions are done, sequence all the different
+ fragments inside gimple_omp_body. */
+
+ new_body = NULL;
+
+ if (ctx->record_type && kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ t = build_fold_addr_expr_loc (loc, ctx->sender_decl);
+ /* fixup_child_record_type might have changed receiver_decl's type. */
+ t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (ctx->receiver_decl, t));
+ }
+
+ if (kind == GF_OMP_TARGET_KIND_REGION)
+ {
+ gimple_seq_add_seq (&new_body, tgt_body);
+ new_body = maybe_catch_exception (new_body);
+ }
+ else if (kind == GF_OMP_TARGET_KIND_DATA)
+ new_body = tgt_body;
+ if (kind != GF_OMP_TARGET_KIND_UPDATE)
+ {
+ gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
+ gimple_omp_set_body (stmt, new_body);
+ }
+
+ bind = gimple_build_bind (NULL, NULL,
+ tgt_bind ? gimple_bind_block (tgt_bind)
+ : NULL_TREE);
+ gsi_replace (gsi_p, bind, true);
+ gimple_bind_add_seq (bind, ilist);
+ gimple_bind_add_stmt (bind, stmt);
+ gimple_bind_add_seq (bind, olist);
+
+ pop_gimplify_context (NULL);
+}
+
+/* Expand code for an OpenMP teams directive. */
+
+static void
+lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+ gimple teams_stmt = gsi_stmt (*gsi_p);
+ struct gimplify_ctx gctx;
+ push_gimplify_context (&gctx);
+
+ tree block = make_node (BLOCK);
+ gimple bind = gimple_build_bind (NULL, NULL, block);
+ gsi_replace (gsi_p, bind, true);
+ gimple_seq bind_body = NULL;
+ gimple_seq dlist = NULL;
+ gimple_seq olist = NULL;
+
+ tree num_teams = find_omp_clause (gimple_omp_teams_clauses (teams_stmt),
+ OMP_CLAUSE_NUM_TEAMS);
+ if (num_teams == NULL_TREE)
+ num_teams = build_int_cst (unsigned_type_node, 0);
+ else
+ {
+ num_teams = OMP_CLAUSE_NUM_TEAMS_EXPR (num_teams);
+ num_teams = fold_convert (unsigned_type_node, num_teams);
+ gimplify_expr (&num_teams, &bind_body, NULL, is_gimple_val, fb_rvalue);
+ }
+ tree thread_limit = find_omp_clause (gimple_omp_teams_clauses (teams_stmt),
+ OMP_CLAUSE_THREAD_LIMIT);
+ if (thread_limit == NULL_TREE)
+ thread_limit = build_int_cst (unsigned_type_node, 0);
+ else
+ {
+ thread_limit = OMP_CLAUSE_THREAD_LIMIT_EXPR (thread_limit);
+ thread_limit = fold_convert (unsigned_type_node, thread_limit);
+ gimplify_expr (&thread_limit, &bind_body, NULL, is_gimple_val,
+ fb_rvalue);
+ }
+
+ lower_rec_input_clauses (gimple_omp_teams_clauses (teams_stmt),
+ &bind_body, &dlist, ctx, NULL);
+ lower_omp (gimple_omp_body_ptr (teams_stmt), ctx);
+ lower_reduction_clauses (gimple_omp_teams_clauses (teams_stmt), &olist, ctx);
+ 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);
+ gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true));
+ gimple_bind_set_body (bind, bind_body);
+
+ pop_gimplify_context (bind);
+
+ gimple_bind_append_vars (bind, ctx->block_vars);
+ BLOCK_VARS (block) = ctx->block_vars;
+ if (BLOCK_VARS (block))
+ TREE_USED (block) = 1;
+}
+
+
/* Callback for lower_omp_1. Return non-NULL if *tp needs to be
regimplified. If DATA is non-NULL, lower_omp_1 is outside
of OpenMP context, but with task_shared_vars set. */
@@ -8760,6 +9597,16 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p
lower_omp_regimplify_p, ctx ? NULL : &wi, NULL))
gimple_regimplify_operands (stmt, gsi_p);
break;
+ case GIMPLE_OMP_TARGET:
+ ctx = maybe_lookup_ctx (stmt);
+ gcc_assert (ctx);
+ lower_omp_target (gsi_p, ctx);
+ break;
+ case GIMPLE_OMP_TEAMS:
+ ctx = maybe_lookup_ctx (stmt);
+ gcc_assert (ctx);
+ lower_omp_teams (gsi_p, ctx);
+ break;
case GIMPLE_CALL:
tree fndecl;
fndecl = gimple_call_fndecl (stmt);
--- gcc/c/c-typeck.c.jj 2013-08-27 20:49:05.000000000 +0200
+++ gcc/c/c-typeck.c 2013-09-04 17:21:17.295112407 +0200
@@ -10773,8 +10773,7 @@ c_finish_omp_cancellation_point (locatio
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one,
- bool &pointer_based_p)
+ bool &maybe_zero_len, unsigned int &first_non_one)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -10801,15 +10800,11 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- if (POINTER_TYPE_P (TREE_TYPE (t))
- && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
- pointer_based_p = true;
return t;
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one,
- pointer_based_p);
+ maybe_zero_len, first_non_one);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -10989,16 +10984,12 @@ handle_omp_array_sections_1 (tree c, tre
}
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
types.safe_push (TREE_TYPE (ret));
- /* For pointer based array sections we will need to evaluate lb more
- than once. */
- if (pointer_based_p)
+ /* We will need to evaluate lb more than once. */
+ tree lb = c_save_expr (low_bound);
+ if (lb != low_bound)
{
- tree lb = c_save_expr (low_bound);
- if (lb != low_bound)
- {
- TREE_PURPOSE (t) = lb;
- low_bound = lb;
- }
+ TREE_PURPOSE (t) = lb;
+ low_bound = lb;
}
ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
return ret;
@@ -11010,12 +11001,10 @@ static bool
handle_omp_array_sections (tree c)
{
bool maybe_zero_len = false;
- bool pointer_based_p = false;
unsigned int first_non_one = 0;
vec<tree> types = vNULL;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one,
- pointer_based_p);
+ maybe_zero_len, first_non_one);
if (first == error_mark_node)
{
types.release ();
@@ -11047,6 +11036,7 @@ handle_omp_array_sections (tree c)
}
if (tem)
first = build2 (COMPOUND_EXPR, TREE_TYPE (first), tem, first);
+ first = c_fully_fold (first, false, NULL);
OMP_CLAUSE_DECL (c) = first;
}
else
@@ -11155,27 +11145,31 @@ handle_omp_array_sections (tree c)
types.release ();
if (side_effects)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
+ first = c_fully_fold (first, false, NULL);
OMP_CLAUSE_DECL (c) = first;
+ if (size)
+ size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
- if (pointer_based_p)
- {
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
- if (!c_mark_addressable (t))
- return false;
- OMP_CLAUSE_DECL (c2) = t;
- t = build_fold_addr_expr (first);
- t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
- build_pointer_type (char_type_node), t);
- t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
- ptrdiff_type_node, t,
- fold_convert_loc (OMP_CLAUSE_LOCATION (c),
- TREE_TYPE (t),
- OMP_CLAUSE_DECL (c2)));
- OMP_CLAUSE_SIZE (c2) = t;
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = c2;
- }
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ return false;
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
+ if (!c_mark_addressable (t))
+ return false;
+ OMP_CLAUSE_DECL (c2) = t;
+ t = build_fold_addr_expr (first);
+ t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
+ tree ptr = OMP_CLAUSE_DECL (c2);
+ if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
+ ptr = build_fold_addr_expr (ptr);
+ t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
+ ptrdiff_type_node, t,
+ fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, ptr));
+ t = c_fully_fold (t, false, NULL);
+ OMP_CLAUSE_SIZE (c2) = t;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
}
return false;
}
@@ -11462,18 +11456,21 @@ c_finish_omp_clauses (tree clauses)
}
else if (!c_mark_addressable (t))
remove = true;
- else if (!COMPLETE_TYPE_P (TREE_TYPE (t)))
+ else if (!COMPLETE_TYPE_P (TREE_TYPE (t))
+ && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD does not have a mappable type in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
- break;
else if (bitmap_bit_p (&generic_head, DECL_UID (t)))
{
- error ("%qD appears more than once in motion clauses", t);
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion clauses", t);
+ else
+ error ("%qD appears more than once in map clauses", t);
remove = true;
}
else
@@ -11509,6 +11506,7 @@ c_finish_omp_clauses (tree clauses)
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
+ case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_ORDERED:
--- gcc/c/c-parser.c.jj 2013-08-27 20:49:04.000000000 +0200
+++ gcc/c/c-parser.c 2013-09-02 11:41:51.023319550 +0200
@@ -12068,7 +12068,10 @@ c_parser_omp_target_data (location_t loc
OMP_TARGET_DATA_CLAUSES (stmt)
= c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data");
- OMP_TARGET_DATA_BODY (stmt) = c_parser_omp_structured_block (parser);
+ keep_next_level ();
+ tree block = c_begin_compound_stmt (true);
+ add_stmt (c_parser_omp_structured_block (parser));
+ OMP_TARGET_DATA_BODY (stmt) = c_end_compound_stmt (loc, block, true);
SET_EXPR_LOCATION (stmt, loc);
return add_stmt (stmt);
@@ -12161,6 +12164,7 @@ c_parser_omp_target (c_parser *parser, e
c_parser_consume_token (parser);
strcpy (p_name, "#pragma omp target");
+ keep_next_level ();
tree block = c_begin_compound_stmt (true);
tree ret = c_parser_omp_teams (loc, parser, p_name,
OMP_TARGET_CLAUSE_MASK, cclauses);
@@ -12182,7 +12186,10 @@ c_parser_omp_target (c_parser *parser, e
OMP_TARGET_CLAUSES (stmt)
= c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
"#pragma omp target");
- OMP_TARGET_BODY (stmt) = c_parser_omp_structured_block (parser);
+ keep_next_level ();
+ tree block = c_begin_compound_stmt (true);
+ add_stmt (c_parser_omp_structured_block (parser));
+ OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
SET_EXPR_LOCATION (stmt, loc);
add_stmt (stmt);
--- gcc/tree.h.jj 2013-08-27 22:05:43.000000000 +0200
+++ gcc/tree.h 2013-09-03 14:35:36.350942556 +0200
@@ -629,6 +629,9 @@ struct GTY(()) tree_base {
OMP_CLAUSE_LINEAR_NO_COPYIN in
OMP_CLAUSE_LINEAR
+ OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION in
+ OMP_CLAUSE_MAP
+
TRANSACTION_EXPR_RELAXED in
TRANSACTION_EXPR
@@ -2050,15 +2053,21 @@ enum omp_clause_map_kind
OMP_CLAUSE_MAP_TO,
OMP_CLAUSE_MAP_FROM,
OMP_CLAUSE_MAP_TOFROM,
- /* This following is an internal only map kind, used for pointer based array
- sections. OMP_CLAUSE_SIZE for these is not the pointer size, which is
- implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */
+ /* The following kind is an internal only map kind, used for pointer based
+ array sections. OMP_CLAUSE_SIZE for these is not the pointer size,
+ which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */
OMP_CLAUSE_MAP_POINTER
};
#define OMP_CLAUSE_MAP_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind)
+/* Nonzero if this map clause is for array (rather than pointer) based array
+ section with zero bias. Both the non-decl OMP_CLAUSE_MAP and
+ correspoidng OMP_CLAUSE_MAP_POINTER clause are marked with this flag. */
+#define OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION(NODE) \
+ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.public_flag)
+
enum omp_clause_proc_bind_kind
{
/* Numbers should match omp_proc_bind_t enum in omp.h. */
--- gcc/gimple.def.jj 2013-08-19 12:07:51.000000000 +0200
+++ gcc/gimple.def 2013-09-02 11:55:57.892684945 +0200
@@ -349,12 +349,18 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "g
CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE)
-/* GIMPLE_OMP_TARGET <BODY, CLAUSES> represents
+/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
#pragma omp target {,data,update}
BODY is the sequence of statements inside the target construct
(NULL for target update).
- CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
-DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_SINGLE)
+ CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
+ CHILD_FN is set when outlining the body of the target region.
+ All the statements in BODY are moved into this newly created
+ function when converting OMP constructs into low-GIMPLE.
+ DATA_ARG is a vec of 3 local variables in the parent function
+ containing data to be mapped to CHILD_FN. This is used to
+ implement the MAP clauses. */
+DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL)
/* GIMPLE_OMP_TEAMS <BODY, CLAUSES> represents #pragma omp teams
BODY is the sequence of statements inside the single section.
--- gcc/fortran/types.def.jj 2013-08-19 12:07:51.000000000 +0200
+++ gcc/fortran/types.def 2013-09-04 20:19:40.880566458 +0200
@@ -120,6 +120,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_VPTR_INT
DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
BT_CONST_VOLATILE_PTR)
DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
+DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
@@ -167,6 +168,8 @@ DEF_FUNCTION_TYPE_5 (BT_FN_BOOL_LONG_LON
BT_PTR_LONG, BT_PTR_LONG)
DEF_FUNCTION_TYPE_5 (BT_FN_VOID_SIZE_VPTR_PTR_PTR_INT, BT_VOID, BT_SIZE,
BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT)
+DEF_FUNCTION_TYPE_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR,
BT_BOOL, BT_LONG, BT_LONG, BT_LONG, BT_LONG,
@@ -203,6 +206,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
BT_ULONGLONG, BT_ULONGLONG,
BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+ BT_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
--- gcc/cp/parser.c.jj 2013-08-27 20:50:58.000000000 +0200
+++ gcc/cp/parser.c 2013-09-04 17:37:20.511076432 +0200
@@ -29451,6 +29451,7 @@ cp_parser_omp_target_data (cp_parser *pa
OMP_TARGET_DATA_CLAUSES (stmt)
= cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data", pragma_tok);
+ keep_next_level (true);
OMP_TARGET_DATA_BODY (stmt) = cp_parser_omp_structured_block (parser);
SET_EXPR_LOCATION (stmt, pragma_tok->location);
@@ -29543,6 +29544,7 @@ cp_parser_omp_target (cp_parser *parser,
cp_lexer_consume_token (parser->lexer);
strcpy (p_name, "#pragma omp target");
+ keep_next_level (true);
tree sb = begin_omp_structured_block ();
unsigned save = cp_parser_begin_omp_structured_block (parser);
tree ret = cp_parser_omp_teams (parser, pragma_tok, p_name,
@@ -29566,6 +29568,7 @@ cp_parser_omp_target (cp_parser *parser,
OMP_TARGET_CLAUSES (stmt)
= cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
"#pragma omp target", pragma_tok);
+ keep_next_level (true);
OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser);
SET_EXPR_LOCATION (stmt, pragma_tok->location);
--- gcc/cp/cp-objcp-common.h.jj 2013-03-20 10:07:19.000000000 +0100
+++ gcc/cp/cp-objcp-common.h 2013-09-04 19:33:40.520583302 +0200
@@ -145,6 +145,8 @@ extern void cp_common_init_ts (void);
#define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause
#undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
#define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
+#undef LANG_HOOKS_OMP_MAPPABLE_TYPE
+#define LANG_HOOKS_OMP_MAPPABLE_TYPE cp_omp_mappable_type
#undef LANG_HOOKS_EH_USE_CXA_END_CLEANUP
#define LANG_HOOKS_EH_USE_CXA_END_CLEANUP true
--- gcc/cp/semantics.c.jj 2013-08-27 20:50:56.000000000 +0200
+++ gcc/cp/semantics.c 2013-09-04 17:21:48.507946784 +0200
@@ -4101,8 +4101,7 @@ cxx_omp_create_clause_info (tree c, tree
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one,
- bool &pointer_based_p)
+ bool &maybe_zero_len, unsigned int &first_non_one)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4134,16 +4133,11 @@ handle_omp_array_sections_1 (tree c, tre
return error_mark_node;
}
t = convert_from_reference (t);
- if (!processing_template_decl
- && POINTER_TYPE_P (TREE_TYPE (t))
- && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
- pointer_based_p = true;
return t;
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one,
- pointer_based_p);
+ maybe_zero_len, first_non_one);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -4326,16 +4320,12 @@ handle_omp_array_sections_1 (tree c, tre
}
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
types.safe_push (TREE_TYPE (ret));
- /* For pointer based array sections we will need to evaluate lb more
- than once. */
- if (pointer_based_p)
+ /* We will need to evaluate lb more than once. */
+ tree lb = cp_save_expr (low_bound);
+ if (lb != low_bound)
{
- tree lb = cp_save_expr (low_bound);
- if (lb != low_bound)
- {
- TREE_PURPOSE (t) = lb;
- low_bound = lb;
- }
+ TREE_PURPOSE (t) = lb;
+ low_bound = lb;
}
ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, false);
return ret;
@@ -4347,12 +4337,10 @@ static bool
handle_omp_array_sections (tree c)
{
bool maybe_zero_len = false;
- bool pointer_based_p = false;
unsigned int first_non_one = 0;
vec<tree> types = vNULL;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one,
- pointer_based_p);
+ maybe_zero_len, first_non_one);
if (first == error_mark_node)
{
types.release ();
@@ -4506,26 +4494,27 @@ handle_omp_array_sections (tree c)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
OMP_CLAUSE_DECL (c) = first;
OMP_CLAUSE_SIZE (c) = size;
- if (pointer_based_p)
- {
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
- if (!cxx_mark_addressable (t))
- return false;
- OMP_CLAUSE_DECL (c2) = t;
- t = build_fold_addr_expr (first);
- t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
- build_pointer_type (char_type_node), t);
- t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
- ptrdiff_type_node, t,
- fold_convert_loc (OMP_CLAUSE_LOCATION (c),
- TREE_TYPE (t),
- OMP_CLAUSE_DECL (c2)));
- OMP_CLAUSE_SIZE (c2) = t;
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = c2;
- }
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ return false;
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
+ if (!cxx_mark_addressable (t))
+ return false;
+ OMP_CLAUSE_DECL (c2) = t;
+ t = build_fold_addr_expr (first);
+ t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, t);
+ tree ptr = OMP_CLAUSE_DECL (c2);
+ if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
+ ptr = build_fold_addr_expr (ptr);
+ t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
+ ptrdiff_type_node, t,
+ fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, ptr));
+ OMP_CLAUSE_SIZE (c2) = t;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
}
}
return false;
@@ -4972,21 +4961,24 @@ finish_omp_clauses (tree clauses)
&& TREE_CODE (TREE_TYPE (t)) != REFERENCE_TYPE
&& !cxx_mark_addressable (t))
remove = true;
- else if (!cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
- == REFERENCE_TYPE)
- ? TREE_TYPE (TREE_TYPE (t))
- : TREE_TYPE (t)))
+ else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
+ == REFERENCE_TYPE)
+ ? TREE_TYPE (TREE_TYPE (t))
+ : TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD does not have a mappable type in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
- break;
else if (bitmap_bit_p (&generic_head, DECL_UID (t)))
{
- error ("%qD appears more than once in motion clauses", t);
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion clauses", t);
+ else
+ error ("%qD appears more than once in map clauses", t);
remove = true;
}
else
--- gcc/langhooks.h.jj 2013-03-20 10:08:13.000000000 +0100
+++ gcc/langhooks.h 2013-09-04 19:29:20.429909391 +0200
@@ -111,6 +111,9 @@ struct lang_hooks_for_types
firstprivate variables. */
void (*omp_firstprivatize_type_sizes) (struct gimplify_omp_ctx *, tree);
+ /* Return true if TYPE is a mappable type. */
+ bool (*omp_mappable_type) (tree type);
+
/* Return TRUE if TYPE1 and TYPE2 are identical for type hashing purposes.
Called only after doing all language independent checks.
At present, this function is only called when both TYPE1 and TYPE2 are
--- gcc/langhooks-def.h.jj 2013-03-20 10:08:41.000000000 +0100
+++ gcc/langhooks-def.h 2013-09-04 19:29:45.450781052 +0200
@@ -77,6 +77,7 @@ extern tree lhd_omp_assignment (tree, tr
struct gimplify_omp_ctx;
extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
tree);
+extern bool lhd_omp_mappable_type (tree);
#define LANG_HOOKS_NAME "GNU unknown"
#define LANG_HOOKS_IDENTIFIER_SIZE sizeof (struct lang_identifier)
@@ -166,6 +167,7 @@ extern tree lhd_make_node (enum tree_cod
#define LANG_HOOKS_TYPE_MAX_SIZE lhd_return_null_const_tree
#define LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES \
lhd_omp_firstprivatize_type_sizes
+#define LANG_HOOKS_OMP_MAPPABLE_TYPE lhd_omp_mappable_type
#define LANG_HOOKS_TYPE_HASH_EQ NULL
#define LANG_HOOKS_GET_ARRAY_DESCR_INFO NULL
#define LANG_HOOKS_GET_SUBRANGE_BOUNDS NULL
@@ -184,6 +186,7 @@ extern tree lhd_make_node (enum tree_cod
LANG_HOOKS_INCOMPLETE_TYPE_ERROR, \
LANG_HOOKS_TYPE_MAX_SIZE, \
LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES, \
+ LANG_HOOKS_OMP_MAPPABLE_TYPE, \
LANG_HOOKS_TYPE_HASH_EQ, \
LANG_HOOKS_GET_ARRAY_DESCR_INFO, \
LANG_HOOKS_GET_SUBRANGE_BOUNDS, \
--- gcc/gimple.h.jj 2013-08-27 21:55:11.000000000 +0200
+++ gcc/gimple.h 2013-09-02 11:57:32.143164024 +0200
@@ -4702,7 +4702,7 @@ static inline tree
gimple_omp_target_clauses (const_gimple gs)
{
GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
- return gs->gimple_omp_single.clauses;
+ return gs->gimple_omp_parallel.clauses;
}
@@ -4712,7 +4712,7 @@ static inline tree *
gimple_omp_target_clauses_ptr (gimple gs)
{
GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
- return &gs->gimple_omp_single.clauses;
+ return &gs->gimple_omp_parallel.clauses;
}
@@ -4722,7 +4722,7 @@ static inline void
gimple_omp_target_set_clauses (gimple gs, tree clauses)
{
GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
- gs->gimple_omp_single.clauses = clauses;
+ gs->gimple_omp_parallel.clauses = clauses;
}
@@ -4747,6 +4747,67 @@ gimple_omp_target_set_kind (gimple g, in
}
+/* Return the child function used to hold the body of OMP_TARGET GS. */
+
+static inline tree
+gimple_omp_target_child_fn (const_gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ return gs->gimple_omp_parallel.child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of
+ OMP_TARGET GS. */
+
+static inline tree *
+gimple_omp_target_child_fn_ptr (gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ return &gs->gimple_omp_parallel.child_fn;
+}
+
+
+/* Set CHILD_FN to be the child function for OMP_TARGET GS. */
+
+static inline void
+gimple_omp_target_set_child_fn (gimple gs, tree child_fn)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ gs->gimple_omp_parallel.child_fn = child_fn;
+}
+
+
+/* Return the artificial argument used to send variables and values
+ from the parent to the children threads in OMP_TARGET GS. */
+
+static inline tree
+gimple_omp_target_data_arg (const_gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ return gs->gimple_omp_parallel.data_arg;
+}
+
+
+/* Return a pointer to the data argument for OMP_TARGET GS. */
+
+static inline tree *
+gimple_omp_target_data_arg_ptr (gimple gs)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ return &gs->gimple_omp_parallel.data_arg;
+}
+
+
+/* Set DATA_ARG to be the data argument for OMP_TARGET GS. */
+
+static inline void
+gimple_omp_target_set_data_arg (gimple gs, tree data_arg)
+{
+ GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+ gs->gimple_omp_parallel.data_arg = data_arg;
+}
+
+
/* Return the clauses associated with OMP_TEAMS GS. */
static inline tree
--- gcc/omp-builtins.def.jj 2013-08-19 12:07:51.000000000 +0200
+++ gcc/omp-builtins.def 2013-09-04 15:14:21.049713033 +0200
@@ -225,3 +225,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_C
BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end",
BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target",
+ BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
+ ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data",
+ BT_FN_VOID, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
+ BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
--- gcc/gimple-pretty-print.c.jj 2013-08-27 21:53:53.000000000 +0200
+++ gcc/gimple-pretty-print.c 2013-08-30 17:53:14.201347484 +0200
@@ -1277,6 +1277,13 @@ dump_gimple_omp_target (pretty_printer *
pp_string (buffer, "#pragma omp target");
pp_string (buffer, kind);
dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
+ if (gimple_omp_target_child_fn (gs))
+ {
+ pp_string (buffer, " [child fn: ");
+ dump_generic_node (buffer, gimple_omp_target_child_fn (gs),
+ spc, flags, false);
+ pp_right_bracket (buffer);
+ }
if (!gimple_seq_empty_p (gimple_omp_body (gs)))
{
newline_and_indent (buffer, spc + 2);
--- gcc/gimplify.c.jj 2013-08-27 22:03:35.000000000 +0200
+++ gcc/gimplify.c 2013-09-04 19:36:07.442834807 +0200
@@ -5976,7 +5976,16 @@ omp_notice_variable (struct gimplify_omp
if (ctx->region_type == ORT_TARGET)
{
if (n == NULL)
- omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ {
+ if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+ {
+ error ("%qD referenced in target region does not have "
+ "a mappable type", decl);
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
+ }
+ else
+ omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ }
else
n->value |= flags;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
@@ -6283,7 +6292,6 @@ gimplify_scan_omp_clauses (tree *list_p,
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
- notice_outer = false;
goto do_add;
case OMP_CLAUSE_DEPEND:
@@ -6500,31 +6508,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
if (private_debug)
code = OMP_CLAUSE_PRIVATE;
else if (flags & GOVD_MAP)
- {
- /* If decl is already in the enclosing device data environment,
- the spec says that it should just be used and no init/assignment
- should be done. If there was any privatization in between though,
- it means that original decl might be in the enclosing device data
- environment, but the privatized might not. */
- struct gimplify_omp_ctx *ctx;
- for (ctx = gimplify_omp_ctxp->outer_context;
- ctx; ctx = ctx->outer_context)
- {
- n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
- if (n == NULL)
- continue;
- if (ctx->region_type == ORT_TARGET_DATA)
- {
- if ((n->value & GOVD_MAP) != 0)
- return 0;
- }
- else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
- | GOVD_PRIVATE | GOVD_REDUCTION
- | GOVD_LINEAR)) != 0)
- break;
- }
- code = OMP_CLAUSE_MAP;
- }
+ code = OMP_CLAUSE_MAP;
else if (flags & GOVD_SHARED)
{
if (is_global_var (decl))
@@ -6689,31 +6673,6 @@ gimplify_adjust_omp_clauses (tree *list_
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
remove = true;
- else
- {
- /* If decl is already in the enclosing device data environment,
- the spec says that it should just be used and no init/assignment
- should be done. If there was any privatization in between though,
- it means that original decl might be in the enclosing device data
- environment, but the privatized might not. */
- struct gimplify_omp_ctx *octx;
- for (octx = ctx->outer_context; octx; octx = octx->outer_context)
- {
- n = splay_tree_lookup (octx->variables,
- (splay_tree_key) decl);
- if (n == NULL)
- continue;
- if (octx->region_type == ORT_TARGET_DATA)
- {
- if ((n->value & GOVD_MAP) != 0)
- remove = true;
- }
- else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
- | GOVD_PRIVATE | GOVD_REDUCTION
- | GOVD_LINEAR)) != 0)
- break;
- }
- }
break;
case OMP_CLAUSE_REDUCTION:
@@ -6722,6 +6681,7 @@ gimplify_adjust_omp_clauses (tree *list_
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_THREAD_LIMIT:
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEVICE:
case OMP_CLAUSE_SCHEDULE:
@@ -7198,7 +7158,28 @@ gimplify_omp_workshare (tree *expr_p, gi
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
- gimplify_and_add (OMP_BODY (expr), &body);
+ if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+ {
+ struct gimplify_ctx gctx;
+ push_gimplify_context (&gctx);
+ gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
+ if (gimple_code (g) == GIMPLE_BIND)
+ pop_gimplify_context (g);
+ else
+ pop_gimplify_context (NULL);
+ if (ort == ORT_TARGET_DATA)
+ {
+ gimple_seq cleanup = NULL;
+ tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+ g = gimple_build_call (fn, 0);
+ gimple_seq_add_stmt (&cleanup, g);
+ g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
+ body = NULL;
+ gimple_seq_add_stmt (&body, g);
+ }
+ }
+ else
+ gimplify_and_add (OMP_BODY (expr), &body);
gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
switch (TREE_CODE (expr))
--- gcc/testsuite/c-c++-common/gomp/map-1.c.jj 2013-07-06 19:18:46.161478820 +0200
+++ gcc/testsuite/c-c++-common/gomp/map-1.c 2013-09-04 17:00:51.174322082 +0200
@@ -33,7 +33,7 @@ foo (int g[3][10], int h[4][8], int i[2]
#pragma omp target map(to: o[2:5]) /* { dg-error "does not have pointer or array type" } */
;
#pragma omp target map(to: a[:][:]) /* { dg-error "array type length expression is not optional" } */
- bar (&a[0][0]);
+ bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" } */
#pragma omp target map(tofrom: b[-1:]) /* { dg-error "negative low bound in array section" } */
bar (b);
#pragma omp target map(tofrom: c[:-3][:]) /* { dg-error "negative length in array section" } */
--- gcc/builtin-types.def.jj 2013-08-19 12:07:51.000000000 +0200
+++ gcc/builtin-types.def 2013-09-03 08:58:48.902658867 +0200
@@ -345,6 +345,7 @@ DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_VPTR_INT
DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_SIZE_CONST_VPTR, BT_BOOL, BT_SIZE,
BT_CONST_VOLATILE_PTR)
DEF_FUNCTION_TYPE_2 (BT_FN_BOOL_INT_BOOL, BT_BOOL, BT_INT, BT_BOOL)
+DEF_FUNCTION_TYPE_2 (BT_FN_VOID_UINT_UINT, BT_VOID, BT_UINT, BT_UINT)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR_PTR, BT_FN_VOID_PTR_PTR)
@@ -472,6 +473,8 @@ DEF_FUNCTION_TYPE_5 (BT_FN_BOOL_VPTR_PTR
DEF_FUNCTION_TYPE_5 (BT_FN_VOID_OMPFN_PTR_UINT_UINT_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT,
BT_UINT)
+DEF_FUNCTION_TYPE_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_6 (BT_FN_INT_STRING_SIZE_INT_SIZE_CONST_STRING_VALIST_ARG,
BT_INT, BT_STRING, BT_SIZE, BT_INT, BT_SIZE,
@@ -511,6 +514,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
BT_ULONGLONG, BT_ULONGLONG,
BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+ BT_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
--- libgomp/libgomp.map.jj 2013-06-12 14:24:42.000000000 +0200
+++ libgomp/libgomp.map 2013-09-04 15:17:42.694659418 +0200
@@ -218,4 +218,9 @@ GOMP_4.0 {
GOMP_parallel;
GOMP_taskgroup_start;
GOMP_taskgroup_end;
+ GOMP_target;
+ GOMP_target_data;
+ GOMP_target_end_data;
+ GOMP_target_update;
+ GOMP_teams;
} GOMP_3.0;
--- libgomp/Makefile.am.jj 2013-03-20 10:02:06.000000000 +0100
+++ libgomp/Makefile.am 2013-09-04 14:59:19.475033836 +0200
@@ -60,7 +60,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_L
libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
- time.c fortran.c affinity.c
+ time.c fortran.c affinity.c target.c
nodist_noinst_HEADERS = libgomp_f.h
nodist_libsubinclude_HEADERS = omp.h
--- libgomp/Makefile.in.jj 2013-03-20 10:02:05.000000000 +0100
+++ libgomp/Makefile.in 2013-09-04 14:59:52.243864494 +0200
@@ -96,7 +96,7 @@ am_libgomp_la_OBJECTS = alloc.lo barrier
error.lo iter.lo iter_ull.lo loop.lo loop_ull.lo ordered.lo \
parallel.lo sections.lo single.lo task.lo team.lo work.lo \
lock.lo mutex.lo proc.lo sem.lo bar.lo ptrlock.lo time.lo \
- fortran.lo affinity.lo
+ fortran.lo affinity.lo target.lo
libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS)
DEFAULT_INCLUDES = -I.@am__isrc@
depcomp = $(SHELL) $(top_srcdir)/../depcomp
@@ -317,7 +317,7 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_L
libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
- time.c fortran.c affinity.c
+ time.c fortran.c affinity.c target.c
nodist_noinst_HEADERS = libgomp_f.h
nodist_libsubinclude_HEADERS = omp.h
@@ -474,6 +474,7 @@ distclean-compile:
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sections.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sem.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/single.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/task.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/team.Plo@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/time.Plo@am__quote@
--- libgomp/libgomp_g.h.jj 2013-04-05 18:15:57.000000000 +0200
+++ libgomp/libgomp_g.h 2013-09-04 15:25:55.982080005 +0200
@@ -199,4 +199,14 @@ extern bool GOMP_single_start (void);
extern void *GOMP_single_copy_start (void);
extern void GOMP_single_copy_end (void *);
+/* target.c */
+
+extern void GOMP_target (int, void (*) (void *), const char *,
+ size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_data (int, size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_end_data (void);
+extern void GOMP_target_update (int, size_t, void **, size_t *,
+ unsigned char *);
+extern void GOMP_teams (unsigned int, unsigned int);
+
#endif /* LIBGOMP_G_H */
--- libgomp/testsuite/libgomp.c++/for-11.C.jj 2013-06-21 08:38:19.000000000 +0200
+++ libgomp/testsuite/libgomp.c++/for-11.C 2013-09-04 15:28:21.542379101 +0200
@@ -75,10 +75,7 @@ int
main ()
{
int err = 0;
-// FIXME: distribute construct must be closely nested
-// in teams region, but we don't handle target expansions
-// yet. Enable when it works.
-// #pragma omp target teams reduction(|:err)
+ #pragma omp target teams reduction(|:err)
{
err |= test_d_normal ();
err |= test_d_ds128_normal ();
--- libgomp/testsuite/libgomp.c++/target-1.C.jj 2013-09-04 18:55:02.294387946 +0200
+++ libgomp/testsuite/libgomp.c++/target-1.C 2013-09-04 18:54:56.430417616 +0200
@@ -0,0 +1 @@
+#include "../libgomp.c/target-1.c"
--- libgomp/testsuite/libgomp.c/for-3.c.jj 2013-07-06 19:44:29.000000000 +0200
+++ libgomp/testsuite/libgomp.c/for-3.c 2013-09-04 15:28:53.986223342 +0200
@@ -77,10 +77,7 @@ int
main ()
{
int err = 0;
-/* FIXME: distribute construct must be closely nested
- in teams region, but we don't handle target expansions
- yet. Enable when it works. */
-/* #pragma omp target teams reduction(|:err) */
+ #pragma omp target teams reduction(|:err)
{
err |= test_d_normal ();
err |= test_d_ds128_normal ();
--- libgomp/testsuite/libgomp.c/target-1.c.jj 2013-09-04 18:54:18.520609855 +0200
+++ libgomp/testsuite/libgomp.c/target-1.c 2013-09-04 18:54:10.000000000 +0200
@@ -0,0 +1,90 @@
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+void
+fn1 (double *x, double *y, int z)
+{
+ int i;
+ for (i = 0; i < z; i++)
+ {
+ x[i] = i & 31;
+ y[i] = (i & 63) - 30;
+ }
+}
+
+#pragma omp declare target
+int tgtv = 6;
+int
+tgt (void)
+{
+ #pragma omp atomic update
+ tgtv++;
+ return 0;
+}
+#pragma omp end declare target
+
+double
+fn2 (int x, int y, int z)
+{
+ double b[1024], c[1024], s = 0;
+ int i, j;
+ fn1 (b, c, x);
+ #pragma omp target data map(to: b)
+ {
+ #pragma omp target map(tofrom: c)
+ #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
+ #pragma omp distribute dist_schedule(static, 4) collapse(1)
+ for (j=0; j < x; j += y)
+ #pragma omp parallel for reduction(+:s)
+ for (i = j; i < j + y; i++)
+ tgt (), s += b[i] * c[i];
+ #pragma omp target update from(b, tgtv)
+ }
+ return s;
+}
+
+double
+fn3 (int x)
+{
+ double b[1024], c[1024], s = 0;
+ int i;
+ fn1 (b, c, x);
+ #pragma omp target map(to: b, c)
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ tgt (), s += b[i] * c[i];
+ return s;
+}
+
+double
+fn4 (int x, double *p)
+{
+ double b[1024], c[1024], d[1024], s = 0;
+ int i;
+ fn1 (b, c, x);
+ fn1 (d + x, p + x, x);
+ #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)])
+ #pragma omp parallel for reduction(+:s)
+ for (i = 0; i < x; i++)
+ s += b[i] * c[i] + d[x + i] + p[x + i];
+ return s;
+}
+
+int
+main ()
+{
+ double a = fn2 (128, 4, 6);
+ int b = tgtv;
+ double c = fn3 (61);
+ #pragma omp target update from(tgtv)
+ int d = tgtv;
+ double e[1024];
+ double f = fn4 (64, e);
+ if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61
+ || f != 8032.0)
+ abort ();
+ return 0;
+}
--- libgomp/target.c.jj 2013-09-04 14:58:30.488287438 +0200
+++ libgomp/target.c 2013-09-04 15:18:31.363404929 +0200
@@ -0,0 +1,83 @@
+/* Copyright (C) 2013 Free Software Foundation, Inc.
+ Contributed by Jakub Jelinek <jakub@redhat.com>.
+
+ This file is part of the GNU OpenMP Library (libgomp).
+
+ Libgomp 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.
+
+ Libgomp 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.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file handles the maintainence of threads in response to team
+ creation and termination. */
+
+#include "libgomp.h"
+#include <stdlib.h>
+#include <string.h>
+
+static int
+resolve_device (int device)
+{
+ return -1;
+}
+
+/* Called when encountering a target directive. If DEVICE
+ is -1, it means use device-var ICV. If it is -2 (or any other value
+ larger than last available hw device, use host fallback.
+ FN is address of host code, FNNAME corresponding name to lookup
+ in the target code. HOSTADDRS, SIZES and KINDS are arrays
+ with MAPNUM entries, with addresses of the host objects,
+ sizes of the host objects (resp. for pointer kind pointer bias
+ and assumed sizeof (void *) size) and kinds. */
+
+void
+GOMP_target (int device, void (*fn) (void *), const char *fnname,
+ size_t mapnum, void **hostaddrs, size_t *sizes,
+ unsigned char *kinds)
+{
+ if (resolve_device (device) == -1)
+ {
+ fn (hostaddrs);
+ return;
+ }
+}
+
+void
+GOMP_target_data (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
+ unsigned char *kinds)
+{
+ if (resolve_device (device) == -1)
+ return;
+}
+
+void
+GOMP_target_end_data (void)
+{
+}
+
+void
+GOMP_target_update (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
+ unsigned char *kinds)
+{
+ if (resolve_device (device) == -1)
+ return;
+}
+
+void
+GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
+{
+}
Jakub