@@ -1,5 +1,20 @@
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA.
+ (is_gimple_omp_oacc_specifically): Handle it.
+ * gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
+ * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise.
+ * omp-low.c (scan_sharing_clauses, scan_omp_target)
+ (expand_omp_target, lower_omp_target, lower_omp_1): Likewise.
+ * gimple.def (GIMPLE_OMP_TARGET): Update comment.
+ * gimple.c (gimple_build_omp_target): Likewise.
+ (gimple_copy): Catch unimplemented case.
+ * tree-inline.c (remap_gimple_stmt): Likewise.
+ * tree-nested.c (convert_nonlocal_reference_stmt)
+ (convert_local_reference_stmt, convert_gimple_call): Likewise.
+ * oacc-builtins.def (BUILT_IN_GOACC_DATA_START)
+ (BUILT_IN_GOACC_DATA_END): New builtins.
+
* omp-low.c (scan_sharing_clauses): Catch unexpected occurrences
of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP.
@@ -1289,6 +1289,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
case GF_OMP_TARGET_KIND_UPDATE:
kind = " update";
break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ kind = " oacc_data";
+ break;
default:
gcc_unreachable ();
}
@@ -1051,7 +1051,8 @@ gimple_build_omp_single (gimple_seq body, tree clauses)
/* Build a GIMPLE_OMP_TARGET statement.
BODY is the sequence of statements that will be executed.
- CLAUSES are any of the OMP target construct's clauses. */
+ KIND is the kind of target region.
+ CLAUSES are any of the construct's clauses. */
gimple
gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
@@ -1747,6 +1748,7 @@ gimple_copy (gimple stmt)
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
copy_omp_body:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
new_seq = gimple_seq_copy (gimple_omp_body (stmt));
gimple_omp_set_body (copy, new_seq);
break;
@@ -360,6 +360,7 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE)
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
+ #pragma acc data
#pragma omp target {,data,update}
BODY is the sequence of statements inside the target construct
(NULL for target update).
@@ -102,6 +102,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_REGION = 0 << 0,
GF_OMP_TARGET_KIND_DATA = 1 << 0,
GF_OMP_TARGET_KIND_UPDATE = 2 << 0,
+ GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -5684,6 +5685,14 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
{
case GIMPLE_OACC_PARALLEL:
return true;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ return true;
+ default:
+ return false;
+ }
default:
return false;
}
@@ -7023,9 +7023,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
return GS_ALL_DONE;
}
-/* Gimplify the gross structure of other OpenMP constructs.
- In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA
- and OMP_TEAMS. */
+/* Gimplify the gross structure of several OpenACC or OpenMP constructs. */
static void
gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -7033,12 +7031,17 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple stmt;
gimple_seq body = NULL;
- enum omp_region_type ort = ORT_WORKSHARE;
+ enum omp_region_type ort;
switch (TREE_CODE (expr))
{
+ case OACC_DATA:
+ ort = (enum omp_region_type) (ORT_TARGET
+ | ORT_TARGET_MAP_FORCE);
+ break;
case OMP_SECTIONS:
case OMP_SINGLE:
+ ort = ORT_WORKSHARE;
break;
case OMP_TARGET:
ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
@@ -7063,9 +7066,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (NULL);
if (!(ort & ORT_TARGET_OFFLOAD))
{
- gimple_seq cleanup = NULL;
- tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+ enum built_in_function end_ix;
+ switch (TREE_CODE (expr))
+ {
+ case OACC_DATA:
+ end_ix = BUILT_IN_GOACC_DATA_END;
+ break;
+ case OMP_TARGET_DATA:
+ end_ix = BUILT_IN_GOMP_TARGET_END_DATA;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ tree fn = builtin_decl_explicit (end_ix);
g = gimple_build_call (fn, 0);
+ gimple_seq cleanup = NULL;
gimple_seq_add_stmt (&cleanup, g);
g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
body = NULL;
@@ -7078,6 +7093,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
+ case OACC_DATA:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
+ OACC_DATA_CLAUSES (expr));
+ break;
case OMP_SECTIONS:
stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
break;
@@ -8047,7 +8066,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
break;
case OACC_KERNELS:
- case OACC_DATA:
case OACC_HOST_DATA:
case OACC_DECLARE:
case OACC_UPDATE:
@@ -8076,6 +8094,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = gimplify_omp_for (expr_p, pre_p);
break;
+ case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
case OMP_TARGET:
@@ -1,7 +1,7 @@
/* This file contains the definitions and documentation for the
OpenACC builtins used in the GNU compiler.
- Copyright (C) 2013 Free Software Foundation, Inc.
+ Copyright (C) 2013-2014 Free Software Foundation, Inc.
Contributed by Thomas Schwinge <thomas@codesourcery.com>.
@@ -29,3 +29,7 @@ along with GCC; see the file COPYING3. If not see
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
+ BT_FN_VOID, ATTR_NOTHROW_LIST)
@@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree c, decl;
bool scan_array_reductions = false;
+ bool offloaded;
+ switch (gimple_code (ctx->stmt))
+ {
+ case GIMPLE_OACC_PARALLEL:
+ offloaded = true;
+ break;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (ctx->stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ offloaded = true;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ offloaded = false;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ break;
+ default:
+ offloaded = false;
+ }
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
@@ -1669,11 +1693,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& 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
+ target regions that are not offloaded; there is nothing to map for
those. */
- if (!gimple_code_is_oacc (ctx->stmt)
- && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
- && !POINTER_TYPE_P (TREE_TYPE (decl)))
+ if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
if (DECL_P (decl))
@@ -1698,9 +1720,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, true, 7, ctx);
else
install_var_field (decl, true, 3, ctx);
- if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL
- || (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_REGION))
+ if (offloaded)
install_var_local (decl, ctx);
}
}
@@ -1824,8 +1844,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
|| (gimple_omp_target_kind (ctx->stmt)
!= GF_OMP_TARGET_KIND_UPDATE));
- if (!gimple_code_is_oacc (ctx->stmt)
- && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+ if (!offloaded)
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
@@ -2340,7 +2359,7 @@ scan_omp_single (gimple stmt, omp_context *outer_ctx)
layout_type (ctx->record_type);
}
-/* Scan an OpenMP target{, data, update} directive. */
+/* Scan a GIMPLE_OMP_TARGET. */
static void
scan_omp_target (gimple stmt, omp_context *outer_ctx)
@@ -2349,6 +2368,12 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
tree name;
int kind = gimple_omp_target_kind (stmt);
+ if (kind == GF_OMP_TARGET_KIND_OACC_DATA)
+ {
+ gcc_assert (taskreg_nesting_level == 0);
+ gcc_assert (target_nesting_level == 0);
+ }
+
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;
@@ -8218,7 +8243,7 @@ expand_omp_atomic (struct omp_region *region)
}
-/* Expand the OpenMP target{, data, update} directive starting at REGION. */
+/* Expand the GIMPLE_OMP_TARGET starting at REGION. */
static void
expand_omp_target (struct omp_region *region)
@@ -8401,12 +8426,23 @@ expand_omp_target (struct omp_region *region)
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;
+ switch (kind)
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ start_ix = BUILT_IN_GOMP_TARGET;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ start_ix = BUILT_IN_GOMP_TARGET_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ start_ix = BUILT_IN_GOACC_DATA_START;
+ break;
+ default:
+ gcc_unreachable ();
+ }
/* By default, the value of DEVICE is -1 (let runtime library choose)
and there is no conditional. */
@@ -8414,10 +8450,12 @@ expand_omp_target (struct omp_region *region)
device = build_int_cst (integer_type_node, -1);
c = find_omp_clause (clauses, OMP_CLAUSE_IF);
+ gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
if (c)
cond = OMP_CLAUSE_IF_EXPR (c);
c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+ gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
if (c)
{
device = OMP_CLAUSE_DEVICE_ID (c);
@@ -8433,6 +8471,7 @@ expand_omp_target (struct omp_region *region)
(cond ? device : -2). */
if (cond)
{
+ gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
cond = gimple_boolify (cond);
basic_block cond_bb, then_bb, else_bb;
@@ -8523,7 +8562,9 @@ expand_omp_target (struct omp_region *region)
gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
gsi_remove (&gsi, true);
}
- if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+ if ((kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
+ && region->exit)
{
gsi = gsi_last_bb (region->exit);
g = gsi_stmt (gsi);
@@ -10277,7 +10318,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
}
-/* Lower the OpenMP target directive in the current statement
+/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
static void
@@ -10298,7 +10339,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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)
+ else if (kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
tgt_body = gimple_omp_body (stmt);
child_fn = ctx->cb.dst_fn;
@@ -10322,6 +10364,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_MAP_TOFROM:
case OMP_CLAUSE_MAP_POINTER:
break;
+ case OMP_CLAUSE_MAP_FORCE_ALLOC:
+ case OMP_CLAUSE_MAP_FORCE_TO:
+ case OMP_CLAUSE_MAP_FORCE_FROM:
+ case OMP_CLAUSE_MAP_FORCE_TOFROM:
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA);
+ break;
default:
gcc_unreachable ();
}
@@ -10330,6 +10381,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -10373,7 +10426,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
lower_omp (&tgt_body, ctx);
target_nesting_level--;
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
lower_omp (&tgt_body, ctx);
if (kind == GF_OMP_TARGET_KIND_REGION)
@@ -10400,9 +10454,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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 tkind_type;
+ int talign_shift;
+ switch (kind)
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ tkind_type = unsigned_char_type_node;
+ talign_shift = 3;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ tkind_type = short_unsigned_type_node;
+ talign_shift = 8;
+ break;
+ default:
+ gcc_unreachable ();
+ }
TREE_VEC_ELT (t, 2)
- = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
- map_cnt),
+ = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
".omp_data_kinds");
DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
@@ -10515,7 +10585,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
- unsigned char tkind = 0;
+ unsigned HOST_WIDE_INT tkind;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_MAP:
@@ -10530,14 +10600,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
gcc_unreachable ();
}
- unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+ gcc_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift));
+ unsigned HOST_WIDE_INT talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
talign = ceil_log2 (talign);
- tkind |= talign << 3;
+ tkind |= talign << talign_shift;
+ gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
- build_int_cst (unsigned_char_type_node,
- tkind));
+ build_int_cstu (tkind_type, tkind));
if (nc && nc != c)
c = nc;
}
@@ -10589,7 +10660,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq_add_seq (&new_body, tgt_body);
new_body = maybe_catch_exception (new_body);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
new_body = tgt_body;
if (kind != GF_OMP_TARGET_KIND_UPDATE)
{
@@ -10810,6 +10882,8 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GIMPLE_OMP_TARGET:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
+ if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA)
+ gcc_assert (!ctx->cancellable);
lower_omp_target (gsi_p, ctx);
break;
case GIMPLE_OMP_TEAMS:
@@ -1397,6 +1397,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
copy = gimple_build_omp_target
(s1, gimple_omp_target_kind (stmt),
@@ -1307,6 +1307,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
save_suppress = info->suppress_expansion;
convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
@@ -1769,6 +1770,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
save_suppress = info->suppress_expansion;
convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
walk_body (convert_local_reference_stmt, convert_local_reference_op,
@@ -2184,6 +2186,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
break;
@@ -1,3 +1,10 @@
+2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+
+ * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start.
+ * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes.
+ * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New
+ functions.
+
2014-02-20 Thomas Schwinge <thomas@codesourcery.com>
* target.c (gomp_load_plugin_for_device): Don't call dlcose if
@@ -233,5 +233,7 @@ OACC_2.0 {
GOACC_2.0 {
global:
+ GOACC_data_end;
+ GOACC_data_start;
GOACC_parallel;
};
@@ -218,5 +218,8 @@ extern void GOMP_teams (unsigned int, unsigned int);
extern void GOACC_parallel (int, void (*) (void *), const void *,
size_t, void **, size_t *, unsigned short *);
+extern void GOACC_data_start (int, const void *,
+ size_t, void **, size_t *, unsigned short *);
+extern void GOACC_data_end (void);
#endif /* LIBGOMP_G_H */
@@ -1,4 +1,4 @@
-/* Copyright (C) 2013 Free Software Foundation, Inc.
+/* Copyright (C) 2013-2014 Free Software Foundation, Inc.
Contributed by Thomas Schwinge <thomas@codesourcery.com>.
@@ -23,7 +23,7 @@
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
-/* This file handles the OpenACC parallel construct. */
+/* This file handles the OpenACC data and parallel constructs. */
#include "libgomp.h"
#include "libgomp_g.h"
@@ -51,3 +51,33 @@ GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
}
GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds_);
}
+
+
+void
+GOACC_data_start (int device, const void *openmp_target, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+ unsigned char kinds_[mapnum];
+ size_t i;
+
+ /* TODO. Eventually, we'll be interpreting all mapping kinds according to
+ the OpenACC semantics; for now we're re-using what is implemented for
+ OpenMP. */
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i];
+ unsigned char align = kinds[i] >> 8;
+ if (kind > 4)
+ gomp_fatal ("memory mapping kind %x for %zd is not yet supported",
+ kind, i);
+
+ kinds_[i] = kind | align << 3;
+ }
+ GOMP_target_data (device, openmp_target, mapnum, hostaddrs, sizes, kinds_);
+}
+
+void
+GOACC_data_end (void)
+{
+ GOMP_target_end_data ();
+}
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA. (is_gimple_omp_oacc_specifically): Handle it. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise. * omp-low.c (scan_sharing_clauses, scan_omp_target) (expand_omp_target, lower_omp_target, lower_omp_1): Likewise. * gimple.def (GIMPLE_OMP_TARGET): Update comment. * gimple.c (gimple_build_omp_target): Likewise. (gimple_copy): Catch unimplemented case. * tree-inline.c (remap_gimple_stmt): Likewise. * tree-nested.c (convert_nonlocal_reference_stmt) (convert_local_reference_stmt, convert_gimple_call): Likewise. * oacc-builtins.def (BUILT_IN_GOACC_DATA_START) (BUILT_IN_GOACC_DATA_END): New builtins. libgomp/ * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New functions. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208016 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 15 ++++++ gcc/gimple-pretty-print.c | 3 ++ gcc/gimple.c | 4 +- gcc/gimple.def | 1 + gcc/gimple.h | 9 ++++ gcc/gimplify.c | 33 +++++++++--- gcc/oacc-builtins.def | 6 ++- gcc/omp-low.c | 132 ++++++++++++++++++++++++++++++++++++---------- gcc/tree-inline.c | 1 + gcc/tree-nested.c | 3 ++ libgomp/ChangeLog.gomp | 7 +++ libgomp/libgomp.map | 2 + libgomp/libgomp_g.h | 3 ++ libgomp/oacc-parallel.c | 34 +++++++++++- 14 files changed, 213 insertions(+), 40 deletions(-)