@@ -69,7 +69,13 @@ enum gimplify_omp_var_data
GOVD_PRIVATE_OUTER_REF = 1024,
GOVD_LINEAR = 2048,
GOVD_ALIGNED = 4096,
+
+ /* Flags for GOVD_MAP. */
+ /* Don't copy back. */
GOVD_MAP_TO_ONLY = 8192,
+ /* Force a specific behavior (or else, a run-time error). */
+ GOVD_MAP_FORCE = 16384,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -86,7 +92,11 @@ enum omp_region_type
ORT_UNTIED_TASK = 5,
ORT_TEAMS = 8,
ORT_TARGET_DATA = 16,
- ORT_TARGET = 32
+ ORT_TARGET = 32,
+
+ /* Flags for ORT_TARGET. */
+ /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
+ ORT_TARGET_MAP_FORCE = 64
};
/* Gimplify hashtable helper. */
@@ -5430,9 +5440,20 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
copy into or out of the context. */
if (!(flags & GOVD_LOCAL))
{
- nflags = flags & GOVD_MAP
- ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT
- : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
+ if (flags & GOVD_MAP)
+ {
+ nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
+#if 0
+ /* Not sure if this is actually needed; haven't found a case
+ where this would change anything; TODO. */
+ if (flags & GOVD_MAP_FORCE)
+ nflags |= OMP_CLAUSE_MAP_FORCE;
+#endif
+ }
+ else if (flags & GOVD_PRIVATE)
+ nflags = GOVD_PRIVATE;
+ else
+ nflags = GOVD_FIRSTPRIVATE;
nflags |= flags & GOVD_SEEN;
t = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (t) == INDIRECT_REF);
@@ -5501,6 +5522,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
for (octx = ctx; octx; octx = octx->outer_context)
if (octx->region_type & ORT_TARGET)
{
+ gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE));
+
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
{
@@ -5562,19 +5585,45 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if (ctx->region_type & ORT_TARGET)
{
+ unsigned map_force;
+ if (ctx->region_type & ORT_TARGET_MAP_FORCE)
+ map_force = GOVD_MAP_FORCE;
+ else
+ map_force = 0;
if (n == NULL)
{
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);
+ omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT | flags);
}
else
- omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags);
}
else
- n->value |= flags;
+ {
+#if 0
+ /* The following fails for:
+
+ int l = 10;
+ float c[l];
+ #pragma acc parallel copy(c[2:4])
+ {
+ #pragma acc parallel
+ {
+ int t = sizeof c;
+ }
+ }
+
+ ..., which we currently don't have to care about (nesting
+ disabled), but eventually will have to; TODO. */
+ if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT))
+ gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force);
+#endif
+
+ n->value |= flags;
+ }
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
goto do_outer;
}
@@ -5858,6 +5907,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add;
case OMP_CLAUSE_MAP:
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ input_location = OMP_CLAUSE_LOCATION (c);
+ /* TODO. */
+ sorry ("data clause not yet implemented");
+ remove = true;
+ break;
+ default:
+ break;
+ }
if (OMP_CLAUSE_SIZE (c)
&& gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
@@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
{
- OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
- ? OMP_CLAUSE_MAP_TO
- : OMP_CLAUSE_MAP_TOFROM;
+ unsigned map_kind;
+ map_kind = (flags & GOVD_MAP_TO_ONLY
+ ? OMP_CLAUSE_MAP_TO
+ : OMP_CLAUSE_MAP_TOFROM);
+ if (flags & GOVD_MAP_FORCE)
+ map_kind |= OMP_CLAUSE_MAP_FORCE;
+ OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind;
+
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple g;
gimple_seq body = NULL;
+ enum omp_region_type ort =
+ (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE);
- gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p,
- ORT_TARGET);
+ gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort);
push_gimplify_context ();
@@ -1064,6 +1064,8 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
|| !splay_tree_lookup (ctx->field_map, (splay_tree_key) var));
gcc_assert ((mask & 2) == 0 || !ctx->sfield_map
|| !splay_tree_lookup (ctx->sfield_map, (splay_tree_key) var));
+ gcc_assert ((mask & 3) == 3
+ || gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
type = TREE_TYPE (var);
if (mask & 4)
@@ -1611,6 +1613,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
case OMP_CLAUSE_MAP:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
@@ -1630,11 +1633,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
{
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* 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
+ if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
@@ -8709,8 +8712,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
- case OMP_CLAUSE_TO:
- case OMP_CLAUSE_FROM:
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -8797,8 +8798,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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))
@@ -8893,12 +8892,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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 ();
}
@@ -10179,6 +10172,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
+#ifdef ENABLE_CHECKING
+ /* First check what we're prepared to handle in the following. */
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case OMP_CLAUSE_MAP_ALLOC:
+ case OMP_CLAUSE_MAP_TO:
+ case OMP_CLAUSE_MAP_FROM:
+ case OMP_CLAUSE_MAP_TOFROM:
+ case OMP_CLAUSE_MAP_POINTER:
+ break;
+ default:
+ gcc_unreachable ();
+ }
+#endif
+ /* FALLTHRU */
+
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
var = OMP_CLAUSE_DECL (c);
@@ -1125,7 +1125,24 @@ enum omp_clause_map_kind
/* 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 = OMP_CLAUSE_MAP_SPECIAL
+ OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL,
+ /* The following are only valid for OpenACC. */
+ /* Flag to force a specific behavior (or else, a run-time error). */
+ OMP_CLAUSE_MAP_FORCE = 1 << 3,
+ /* Allocate. */
+ OMP_CLAUSE_MAP_FORCE_ALLOC = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_ALLOC,
+ /* ..., and copy to device. */
+ OMP_CLAUSE_MAP_FORCE_TO = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TO,
+ /* ..., and copy from device. */
+ OMP_CLAUSE_MAP_FORCE_FROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_FROM,
+ /* ..., and copy to and from device. */
+ OMP_CLAUSE_MAP_FORCE_TOFROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TOFROM,
+ /* Must already be present. */
+ OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL,
+ /* Deallocate a mapping, without copying from device. */
+ OMP_CLAUSE_MAP_FORCE_DEALLOC,
+ /* Is a device pointer. */
+ OMP_CLAUSE_MAP_FORCE_DEVICEPTR
};
enum omp_clause_proc_bind_kind
@@ -506,6 +506,27 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE_MAP_TOFROM:
pp_string (buffer, "tofrom");
break;
+ case OMP_CLAUSE_MAP_FORCE_ALLOC:
+ pp_string (buffer, "force_alloc");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_TO:
+ pp_string (buffer, "force_to");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_FROM:
+ pp_string (buffer, "force_from");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_TOFROM:
+ pp_string (buffer, "force_tofrom");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ pp_string (buffer, "force_present");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ pp_string (buffer, "force_dealloc");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ pp_string (buffer, "force_deviceptr");
+ break;
default:
gcc_unreachable ();
}
From: Thomas Schwinge <thomas@codesourcery.com> gcc/ * tree-core.h (omp_clause_map_kind): Add OMP_CLAUSE_MAP_FORCE, OMP_CLAUSE_MAP_FORCE_ALLOC, OMP_CLAUSE_MAP_FORCE_TO, OMP_CLAUSE_MAP_FORCE_FROM, OMP_CLAUSE_MAP_FORCE_TOFROM, OMP_CLAUSE_MAP_FORCE_PRESENT, OMP_CLAUSE_MAP_FORCE_DEALLOC, and OMP_CLAUSE_MAP_FORCE_DEVICEPTR. * tree-pretty-print.c (dump_omp_clause): Handle these. * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_FORCE. (omp_region_type): Add ORT_TARGET_MAP_FORCE. (omp_add_variable, omp_notice_threadprivate_variable) (omp_notice_variable, gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses_1): Extend accordingly. (gimplify_oacc_parallel): Add ORT_TARGET_MAP_FORCE to ORT_TARGET usage. * omp-low.c (install_var_field, scan_sharing_clauses) (lower_oacc_parallel, lower_omp_target): Extend accordingly. --- gcc/gimplify.c | 92 ++++++++++++++++++++++++++++++++++++++++++------- gcc/omp-low.c | 33 +++++++++++------- gcc/tree-core.h | 19 +++++++++- gcc/tree-pretty-print.c | 21 +++++++++++ 4 files changed, 140 insertions(+), 25 deletions(-)