@@ -1341,6 +1341,11 @@ extern void protected_set_expr_location
#define OMP_TEAMS_COMBINED(NODE) \
(OMP_TEAMS_CHECK (NODE)->base.private_flag)
+/* True on an OMP_TARGET statement if it represents explicit
+ combined target teams, target parallel or target simd constructs. */
+#define OMP_TARGET_COMBINED(NODE) \
+ (OMP_TARGET_CHECK (NODE)->base.private_flag)
+
/* True if OMP_ATOMIC* is supposed to be sequentially consistent
as opposed to relaxed. */
#define OMP_ATOMIC_SEQ_CST(NODE) \
@@ -1445,13 +1450,21 @@ extern void protected_set_expr_location
((enum gomp_map_kind) OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind)
#define OMP_CLAUSE_SET_MAP_KIND(NODE, MAP_KIND) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \
- = (unsigned char) (MAP_KIND))
+ = (unsigned int) (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 corresponding
OMP_CLAUSE_MAP with GOMP_MAP_POINTER 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)
+/* Nonzero if the same decl appears both in OMP_CLAUSE_MAP and either
+ OMP_CLAUSE_PRIVATE or OMP_CLAUSE_FIRSTPRIVATE. */
+#define OMP_CLAUSE_MAP_PRIVATE(NODE) \
+ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero if this is a mapped array section, that might need special
+ treatment if OMP_CLAUSE_SIZE is zero. */
+#define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \
+ TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
#define OMP_CLAUSE_PROC_BIND_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)
@@ -1354,7 +1354,7 @@ struct GTY(()) tree_omp_clause {
enum omp_clause_schedule_kind schedule_kind;
enum omp_clause_depend_kind depend_kind;
/* See include/gomp-constants.h for enum gomp_map_kind's values. */
- unsigned char map_kind;
+ unsigned int map_kind;
enum omp_clause_proc_bind_kind proc_bind_kind;
enum tree_code reduction_code;
enum omp_clause_linear_kind linear_kind;
@@ -90,6 +90,8 @@ enum gimplify_omp_var_data
/* Flag for GOVD_LINEAR or GOVD_LASTPRIVATE: no outer reference. */
GOVD_LINEAR_LASTPRIVATE_NO_OUTER = 16384,
+ GOVD_MAP_0LEN_ARRAY = 32768,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -110,6 +112,7 @@ enum omp_region_type
ORT_TARGET_DATA = 16,
/* Data region with offloading. */
ORT_TARGET = 32,
+ ORT_COMBINED_TARGET = 33,
/* Dummy OpenMP region, used to disable expansion of
DECL_VALUE_EXPRs in taskloop pre body. */
ORT_NONE = 64
@@ -156,6 +159,9 @@ struct gimplify_omp_ctx
enum omp_region_type region_type;
bool combined_loop;
bool distribute;
+ bool target_map_scalars_firstprivate;
+ bool target_map_pointers_as_0len_arrays;
+ bool target_firstprivatize_array_bases;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -2260,7 +2266,7 @@ maybe_fold_stmt (gimple_stmt_iterator *g
{
struct gimplify_omp_ctx *ctx;
for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
- if (ctx->region_type == ORT_TARGET)
+ if ((ctx->region_type & ORT_TARGET) != 0)
return false;
return fold_stmt (gsi);
}
@@ -5561,8 +5567,13 @@ omp_firstprivatize_variable (struct gimp
else
return;
}
- else if (ctx->region_type == ORT_TARGET)
- omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+ else if ((ctx->region_type & ORT_TARGET) != 0)
+ {
+ if (ctx->target_map_scalars_firstprivate)
+ omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
+ else
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+ }
else if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD
&& ctx->region_type != ORT_TARGET_DATA)
@@ -5648,7 +5659,7 @@ omp_add_variable (struct gimplify_omp_ct
flags |= GOVD_SEEN;
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- if (n != NULL && n->value != GOVD_ALIGNED)
+ if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
{
/* We shouldn't be re-adding the decl with the same data
sharing class. */
@@ -5678,6 +5689,9 @@ omp_add_variable (struct gimplify_omp_ct
nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
else if (flags & GOVD_PRIVATE)
nflags = GOVD_PRIVATE;
+ else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+ && (flags & GOVD_FIRSTPRIVATE))
+ nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
else
nflags = GOVD_FIRSTPRIVATE;
nflags |= flags & GOVD_SEEN;
@@ -5746,7 +5760,7 @@ omp_notice_threadprivate_variable (struc
struct gimplify_omp_ctx *octx;
for (octx = ctx; octx; octx = octx->outer_context)
- if (octx->region_type == ORT_TARGET)
+ if ((octx->region_type & ORT_TARGET) != 0)
{
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
@@ -5810,19 +5824,66 @@ omp_notice_variable (struct gimplify_omp
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- if (ctx->region_type == ORT_TARGET)
+ if ((ctx->region_type & ORT_TARGET) != 0)
{
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
if (n == NULL)
{
- if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+ unsigned nflags = flags;
+ if (ctx->target_map_pointers_as_0len_arrays
+ || ctx->target_map_scalars_firstprivate)
+ {
+ bool is_declare_target = false;
+ bool is_scalar = false;
+ if (is_global_var (decl)
+ && varpool_node::get_create (decl)->offloadable)
+ {
+ 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
+ && (n->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED
+ && (n->value & GOVD_DATA_SHARE_CLASS) != 0)
+ break;
+ }
+ is_declare_target = octx == NULL;
+ }
+ if (!is_declare_target && ctx->target_map_scalars_firstprivate)
+ {
+ tree type = TREE_TYPE (decl);
+ if (TREE_CODE (type) == REFERENCE_TYPE)
+ type = TREE_TYPE (type);
+ if (TREE_CODE (type) == COMPLEX_TYPE)
+ type = TREE_TYPE (type);
+ if (INTEGRAL_TYPE_P (type)
+ || SCALAR_FLOAT_TYPE_P (type)
+ || TREE_CODE (type) == POINTER_TYPE)
+ is_scalar = true;
+ }
+ if (is_declare_target)
+ ;
+ else if (ctx->target_map_pointers_as_0len_arrays
+ && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+ || (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (TREE_TYPE (decl)))
+ == POINTER_TYPE)))
+ nflags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY;
+ else if (is_scalar)
+ nflags |= GOVD_FIRSTPRIVATE;
+ }
+ if (nflags == flags
+ && !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);
+ nflags |= GOVD_MAP | GOVD_EXPLICIT;
}
- else
- omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ else if (nflags == flags)
+ nflags |= GOVD_MAP;
+ omp_add_variable (ctx, decl, nflags);
}
else
{
@@ -6144,6 +6205,24 @@ gimplify_scan_omp_clauses (tree *list_p,
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
+ if (code == OMP_TARGET && !lang_GNU_Fortran ())
+ {
+ ctx->target_map_pointers_as_0len_arrays = true;
+ /* FIXME: For Fortran we want to set this too, when
+ the Fortran FE is updated to OpenMP 4.1. */
+ ctx->target_map_scalars_firstprivate = true;
+ }
+ if (!lang_GNU_Fortran ())
+ switch (code)
+ {
+ case OMP_TARGET:
+ case OMP_TARGET_DATA:
+ case OMP_TARGET_ENTER_DATA:
+ case OMP_TARGET_EXIT_DATA:
+ ctx->target_firstprivatize_array_bases = true;
+ default:
+ break;
+ }
while ((c = *list_p) != NULL)
{
@@ -6290,11 +6369,18 @@ gimplify_scan_omp_clauses (tree *list_p,
&& ctx->region_type == ORT_WORKSHARE
&& octx == outer_ctx)
flags = GOVD_SEEN | GOVD_SHARED;
+ else if (octx
+ && octx->region_type == ORT_COMBINED_TARGET)
+ flags &= ~GOVD_LASTPRIVATE;
else
break;
- gcc_checking_assert (splay_tree_lookup (octx->variables,
- (splay_tree_key)
- decl) == NULL);
+ splay_tree_node on
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ gcc_assert (on == NULL
+ || (octx->region_type == ORT_COMBINED_TARGET
+ && (on->value
+ & GOVD_DATA_SHARE_CLASS) == 0));
omp_add_variable (octx, decl, flags);
if (octx->outer_context == NULL)
break;
@@ -6319,10 +6405,24 @@ gimplify_scan_omp_clauses (tree *list_p,
case OMP_CLAUSE_MAP:
decl = OMP_CLAUSE_DECL (c);
if (error_operand_p (decl))
+ remove = true;
+ switch (code)
{
- remove = true;
+ case OMP_TARGET:
+ break;
+ case OMP_TARGET_DATA:
+ case OMP_TARGET_ENTER_DATA:
+ case OMP_TARGET_EXIT_DATA:
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ /* For target {,enter ,exit }data only the array slice is
+ mapped, but not the pointer to it. */
+ remove = true;
+ break;
+ default:
break;
}
+ if (remove)
+ break;
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -6332,6 +6432,14 @@ gimplify_scan_omp_clauses (tree *list_p,
remove = true;
break;
}
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+ {
+ OMP_CLAUSE_SIZE (c)
+ = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL);
+ omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+ GOVD_FIRSTPRIVATE | GOVD_SEEN);
+ }
if (!DECL_P (decl))
{
if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
@@ -6643,7 +6751,10 @@ gimplify_scan_omp_clauses (tree *list_p,
case OMP_CLAUSE_NOGROUP:
case OMP_CLAUSE_THREADS:
case OMP_CLAUSE_SIMD:
+ break;
+
case OMP_CLAUSE_DEFAULTMAP:
+ ctx->target_map_scalars_firstprivate = false;
break;
case OMP_CLAUSE_ALIGNED:
@@ -6759,6 +6870,30 @@ gimplify_adjust_omp_clauses_1 (splay_tre
OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
+ else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
+ {
+ tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == POINTER_TYPE)
+ OMP_CLAUSE_DECL (clause)
+ = build_simple_mem_ref_loc (input_location, decl);
+ OMP_CLAUSE_DECL (clause)
+ = build2 (MEM_REF, char_type_node, OMP_CLAUSE_DECL (clause),
+ build_int_cst (build_pointer_type (char_type_node), 0));
+ OMP_CLAUSE_SIZE (clause) = size_zero_node;
+ OMP_CLAUSE_SIZE (nc) = size_zero_node;
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1;
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_CHAIN (nc) = *list_p;
+ OMP_CLAUSE_CHAIN (clause) = nc;
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0),
+ pre_p, NULL, is_gimple_val, fb_rvalue);
+ gimplify_omp_ctxp = ctx;
+ }
else if (code == OMP_CLAUSE_MAP)
{
OMP_CLAUSE_SET_MAP_KIND (clause,
@@ -6785,7 +6920,10 @@ gimplify_adjust_omp_clauses_1 (splay_tre
OMP_CLAUSE_MAP);
OMP_CLAUSE_DECL (nc) = decl;
OMP_CLAUSE_SIZE (nc) = size_zero_node;
- OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+ if (gimplify_omp_ctxp->target_firstprivatize_array_bases)
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
OMP_CLAUSE_CHAIN (clause) = nc;
}
@@ -6910,12 +7048,14 @@ gimplify_adjust_omp_clauses (gimple_seq
if (!DECL_P (decl))
break;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
- if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN)
+ if ((ctx->region_type & ORT_TARGET) != 0
+ && !(n->value & GOVD_SEEN)
&& !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)
{
/* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because
for these, TREE_CODE (DECL_SIZE (decl)) will always be
@@ -6935,17 +7075,33 @@ gimplify_adjust_omp_clauses (gimple_seq
omp_notice_variable (ctx->outer_context,
OMP_CLAUSE_SIZE (c), true);
}
- tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_DECL (nc) = decl;
- OMP_CLAUSE_SIZE (nc) = size_zero_node;
- OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
- OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = nc;
- c = nc;
+ if (((ctx->region_type & ORT_TARGET) != 0
+ || !ctx->target_firstprivatize_array_bases)
+ && ((n->value & GOVD_SEEN) == 0
+ || (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0))
+ {
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ OMP_CLAUSE_SIZE (nc) = size_zero_node;
+ if (ctx->target_firstprivatize_array_bases)
+ OMP_CLAUSE_SET_MAP_KIND (nc,
+ GOMP_MAP_FIRSTPRIVATE_POINTER);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = nc;
+ c = nc;
+ }
+ }
+ else
+ {
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
+ if ((n->value & GOVD_SEEN)
+ && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)))
+ OMP_CLAUSE_MAP_PRIVATE (c) = 1;
}
- else if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
- OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
break;
case OMP_CLAUSE_TO:
@@ -7888,9 +8044,11 @@ gimplify_omp_workshare (tree *expr_p, gi
case OMP_SINGLE:
ort = ORT_WORKSHARE;
break;
+ case OMP_TARGET:
+ ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
+ break;
case OACC_KERNELS:
case OACC_PARALLEL:
- case OMP_TARGET:
ort = ORT_TARGET;
break;
case OACC_DATA:
@@ -7905,7 +8063,7 @@ gimplify_omp_workshare (tree *expr_p, gi
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
TREE_CODE (expr));
- if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+ if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
{
push_gimplify_context ();
gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
@@ -1071,24 +1071,35 @@ lookup_field (tree var, omp_context *ctx
}
static inline tree
-lookup_sfield (tree var, omp_context *ctx)
+lookup_sfield (splay_tree_key key, omp_context *ctx)
{
splay_tree_node n;
n = splay_tree_lookup (ctx->sfield_map
- ? ctx->sfield_map : ctx->field_map,
- (splay_tree_key) var);
+ ? ctx->sfield_map : ctx->field_map, key);
return (tree) n->value;
}
static inline tree
-maybe_lookup_field (tree var, omp_context *ctx)
+lookup_sfield (tree var, omp_context *ctx)
+{
+ return lookup_sfield ((splay_tree_key) var, ctx);
+}
+
+static inline tree
+maybe_lookup_field (splay_tree_key key, omp_context *ctx)
{
splay_tree_node n;
- n = splay_tree_lookup (ctx->field_map, (splay_tree_key) var);
+ n = splay_tree_lookup (ctx->field_map, key);
return n ? (tree) n->value : NULL_TREE;
}
static inline tree
+maybe_lookup_field (tree var, omp_context *ctx)
+{
+ return maybe_lookup_field ((splay_tree_key) var, ctx);
+}
+
+static inline tree
lookup_oacc_reduction (const char *id, omp_context *ctx)
{
splay_tree_node n;
@@ -1359,12 +1370,18 @@ build_outer_var_ref (tree var, omp_conte
/* Build tree nodes to access the field for VAR on the sender side. */
static tree
-build_sender_ref (tree var, omp_context *ctx)
+build_sender_ref (splay_tree_key key, omp_context *ctx)
{
- tree field = lookup_sfield (var, ctx);
+ tree field = lookup_sfield (key, ctx);
return omp_build_component_ref (ctx->sender_decl, field);
}
+static tree
+build_sender_ref (tree var, omp_context *ctx)
+{
+ return build_sender_ref ((splay_tree_key) var, ctx);
+}
+
/* Add a new field for VAR inside the structure CTX->SENDER_DECL. */
static void
@@ -1908,6 +1925,17 @@ scan_sharing_clauses (tree clauses, omp_
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
do_private:
+ if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ && is_gimple_omp_offloaded (ctx->stmt))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ install_var_field (decl, !is_reference (decl), 3, ctx);
+ else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ install_var_field (decl, true, 3, ctx);
+ else
+ install_var_field (decl, false, 3, ctx);
+ }
if (is_variable_sized (decl))
{
if (is_task_ctx (ctx))
@@ -1930,10 +1958,6 @@ scan_sharing_clauses (tree clauses, omp_
else if (!global)
install_var_field (decl, by_ref, 3, ctx);
}
- else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
- || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
- && is_gimple_omp_offloaded (ctx->stmt))
- install_var_field (decl, !is_reference (decl), 3, ctx);
install_var_local (decl, ctx);
if (is_gimple_omp_oacc (ctx->stmt)
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
@@ -1944,9 +1968,9 @@ scan_sharing_clauses (tree clauses, omp_
tree ptype = build_pointer_type (type);
tree array = create_tmp_var (ptype,
oacc_get_reduction_array_id (var));
- omp_context *c = (ctx->field_map ? ctx : ctx->outer);
- install_var_field (array, true, 3, c);
- install_var_local (array, c);
+ omp_context *octx = (ctx->field_map ? ctx : ctx->outer);
+ install_var_field (array, true, 3, octx);
+ install_var_local (array, octx);
/* Insert it into the current context. */
splay_tree_insert (ctx->reduction_map, (splay_tree_key)
@@ -1959,6 +1983,23 @@ scan_sharing_clauses (tree clauses, omp_
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
+ decl = OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ install_var_field (decl, true, 3, ctx);
+ else
+ install_var_field (decl, false, 3, ctx);
+ if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ install_var_local (decl2, ctx);
+ }
+ install_var_local (decl, ctx);
+ break;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
goto do_private;
@@ -2025,6 +2066,21 @@ scan_sharing_clauses (tree clauses, omp_
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
break;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (DECL_SIZE (decl)
+ && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ install_var_local (decl2, ctx);
+ }
+ install_var_local (decl, ctx);
+ break;
+ }
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
@@ -2034,7 +2090,11 @@ scan_sharing_clauses (tree clauses, omp_
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
gcc_assert (DECL_P (decl2));
- install_var_field (decl2, true, 3, ctx);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_PRIVATE (c))
+ install_var_field (decl2, true, 11, ctx);
+ else
+ install_var_field (decl2, true, 3, ctx);
install_var_local (decl2, ctx);
install_var_local (decl, ctx);
}
@@ -2045,6 +2105,9 @@ scan_sharing_clauses (tree clauses, omp_
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 7, ctx);
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_PRIVATE (c))
+ install_var_field (decl, true, 11, ctx);
else
install_var_field (decl, true, 3, ctx);
if (is_gimple_omp_offloaded (ctx->stmt))
@@ -2147,11 +2210,23 @@ scan_sharing_clauses (tree clauses, omp_
/* FALLTHRU */
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_LINEAR:
- case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
if (is_variable_sized (decl))
- install_var_local (decl, ctx);
+ {
+ if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ && is_gimple_omp_offloaded (ctx->stmt))
+ {
+ tree decl2 = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ install_var_local (decl2, ctx);
+ fixup_remapped_decl (decl2, ctx, false);
+ }
+ install_var_local (decl, ctx);
+ }
fixup_remapped_decl (decl, ctx,
OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
&& OMP_CLAUSE_PRIVATE_DEBUG (c));
@@ -2201,7 +2276,8 @@ scan_sharing_clauses (tree clauses, omp_
break;
if (DECL_P (decl))
{
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
&& !COMPLETE_TYPE_P (TREE_TYPE (decl)))
{
@@ -2255,6 +2331,7 @@ scan_sharing_clauses (tree clauses, omp_
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_NOGROUP:
case OMP_CLAUSE_DEFAULTMAP:
+ case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE__CILK_FOR_COUNT_:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
@@ -3924,11 +4001,8 @@ handle_simd_reference (location_t loc, t
tree z = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_vard)));
if (TREE_CONSTANT (z))
{
- const char *name = NULL;
- if (DECL_NAME (new_vard))
- name = IDENTIFIER_POINTER (DECL_NAME (new_vard));
-
- z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)), name);
+ z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)),
+ get_name (new_vard));
gimple_add_tmp_var (z);
TREE_ADDRESSABLE (z) = 1;
z = build_fold_addr_expr_loc (loc, z);
@@ -4127,9 +4201,7 @@ lower_rec_input_clauses (tree clauses, g
tree type = TREE_TYPE (d);
gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
- const char *name = NULL;
- if (DECL_NAME (orig_var))
- name = IDENTIFIER_POINTER (DECL_NAME (orig_var));
+ const char *name = get_name (orig_var);
if (TREE_CONSTANT (v))
{
x = create_tmp_var_raw (type, name);
@@ -4139,7 +4211,8 @@ lower_rec_input_clauses (tree clauses, g
}
else
{
- tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
+ tree atmp
+ = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
tree t = maybe_lookup_decl (v, ctx);
if (t)
v = t;
@@ -4152,7 +4225,8 @@ lower_rec_input_clauses (tree clauses, g
t = fold_build2_loc (clause_loc, MULT_EXPR,
TREE_TYPE (v), t,
TYPE_SIZE_UNIT (TREE_TYPE (type)));
- x = build_call_expr_loc (clause_loc, atmp, 1, t);
+ tree al = size_int (TYPE_ALIGN (TREE_TYPE (type)));
+ x = build_call_expr_loc (clause_loc, atmp, 2, t, al);
}
tree ptype = build_pointer_type (TREE_TYPE (type));
@@ -4362,8 +4436,9 @@ lower_rec_input_clauses (tree clauses, g
x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
/* void *tmp = __builtin_alloca */
- atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
- stmt = gimple_build_call (atmp, 1, x);
+ atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+ stmt = gimple_build_call (atmp, 2, x,
+ size_int (DECL_ALIGN (var)));
tmp = create_tmp_var_raw (ptr_type_node);
gimple_add_tmp_var (tmp);
gimple_call_set_lhs (stmt, tmp);
@@ -4400,12 +4475,8 @@ lower_rec_input_clauses (tree clauses, g
x = NULL_TREE;
else
{
- const char *name = NULL;
- if (DECL_NAME (var))
- name = IDENTIFIER_POINTER (DECL_NAME (new_var));
-
x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
- name);
+ get_name (var));
gimple_add_tmp_var (x);
TREE_ADDRESSABLE (x) = 1;
x = build_fold_addr_expr_loc (clause_loc, x);
@@ -4413,8 +4484,11 @@ lower_rec_input_clauses (tree clauses, g
}
else
{
- tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
- x = build_call_expr_loc (clause_loc, atmp, 1, x);
+ tree atmp
+ = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+ tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+ tree al = size_int (TYPE_ALIGN (rtype));
+ x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
}
if (x)
@@ -5489,11 +5563,7 @@ lower_send_clauses (tree clauses, gimple
/* Handle taskloop firstprivate/lastprivate, where the
lastprivate on GIMPLE_OMP_TASK is represented as
OMP_CLAUSE_SHARED_FIRSTPRIVATE. */
- tree f
- = (tree)
- splay_tree_lookup (ctx->sfield_map
- ? ctx->sfield_map : ctx->field_map,
- (splay_tree_key) &DECL_UID (val))->value;
+ tree f = lookup_sfield ((splay_tree_key) &DECL_UID (val), ctx);
x = omp_build_component_ref (ctx->sender_decl, f);
if (use_pointer_for_field (val, ctx))
var = build_fold_addr_expr (var);
@@ -12883,6 +12953,7 @@ lower_omp_target (gimple_stmt_iterator *
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
break;
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
@@ -12918,6 +12989,28 @@ lower_omp_target (gimple_stmt_iterator *
var = var2;
}
+ if (offloaded
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ {
+ tree type = build_pointer_type (TREE_TYPE (var));
+ tree new_var = lookup_decl (var, ctx);
+ x = create_tmp_var_raw (type, get_name (new_var));
+ gimple_add_tmp_var (x);
+ x = build_simple_mem_ref (x);
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ continue;
+ }
+
+ if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c))
+ {
+ map_cnt++;
+ continue;
+ }
+
if (!maybe_lookup_field (var, ctx))
continue;
@@ -12925,6 +13018,7 @@ lower_omp_target (gimple_stmt_iterator *
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
+
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
@@ -12936,14 +13030,70 @@ lower_omp_target (gimple_stmt_iterator *
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_IS_DEVICE_PTR:
map_cnt++;
var = OMP_CLAUSE_DECL (c);
if (!is_reference (var)
&& !is_gimple_reg_type (TREE_TYPE (var)))
{
- x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
+ if (is_variable_sized (var))
+ {
+ tree pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ gcc_assert (DECL_P (pvar));
+ tree new_pvar = lookup_decl (pvar, ctx);
+ x = build_fold_indirect_ref (new_pvar);
+ TREE_THIS_NOTRAP (x) = 1;
+ }
+ else
+ x = build_receiver_ref (var, true, ctx);
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ break;
+
+ case OMP_CLAUSE_PRIVATE:
+ var = OMP_CLAUSE_DECL (c);
+ if (is_variable_sized (var))
+ {
+ tree new_var = lookup_decl (var, ctx);
+ tree pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ gcc_assert (DECL_P (pvar));
+ tree new_pvar = lookup_decl (pvar, ctx);
+ x = build_fold_indirect_ref (new_pvar);
+ TREE_THIS_NOTRAP (x) = 1;
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ break;
+
+ case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_IS_DEVICE_PTR:
+ var = OMP_CLAUSE_DECL (c);
+ map_cnt++;
+ if (is_variable_sized (var))
+ {
+ tree new_var = lookup_decl (var, ctx);
+ tree pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ gcc_assert (DECL_P (pvar));
+ tree new_pvar = lookup_decl (pvar, ctx);
+ x = build_fold_indirect_ref (new_pvar);
+ TREE_THIS_NOTRAP (x) = 1;
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ {
+ tree new_var = lookup_decl (var, ctx);
+ tree type = build_pointer_type (TREE_TYPE (var));
+ x = create_tmp_var_raw (type, get_name (new_var));
+ gimple_add_tmp_var (x);
+ x = build_simple_mem_ref (x);
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
@@ -13013,7 +13163,7 @@ lower_omp_target (gimple_stmt_iterator *
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
- tree ovar, nc, s, purpose, var, x;
+ tree ovar, nc, s, purpose, var, x, type;
unsigned int talign;
default:
@@ -13044,6 +13194,10 @@ lower_omp_target (gimple_stmt_iterator *
}
else
{
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ break;
if (DECL_SIZE (ovar)
&& TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
{
@@ -13053,7 +13207,14 @@ lower_omp_target (gimple_stmt_iterator *
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
- if (!maybe_lookup_field (ovar, ctx))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_PRIVATE (c))
+ {
+ if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar),
+ ctx))
+ continue;
+ }
+ else if (!maybe_lookup_field (ovar, ctx))
continue;
}
@@ -13063,7 +13224,12 @@ lower_omp_target (gimple_stmt_iterator *
if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
- x = build_sender_ref (ovar, ctx);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_PRIVATE (c))
+ x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar),
+ ctx);
+ else
+ x = build_sender_ref (ovar, ctx);
if (maybe_lookup_oacc_reduction (var, ctx))
{
gcc_checking_assert (offloaded
@@ -13101,7 +13267,7 @@ lower_omp_target (gimple_stmt_iterator *
|| map_kind == GOMP_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
- x = build_sender_ref (ovar, ctx);
+ x = unshare_expr (x);
x = build_simple_mem_ref (x);
gimplify_assign (var, x, &olist);
}
@@ -13121,35 +13287,74 @@ lower_omp_target (gimple_stmt_iterator *
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
- unsigned HOST_WIDE_INT tkind;
+ unsigned HOST_WIDE_INT tkind, tkind_zero;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_MAP:
tkind = OMP_CLAUSE_MAP_KIND (c);
+ tkind_zero = tkind;
+ if (OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c))
+ switch (tkind)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION;
+ break;
+ default:
+ break;
+ }
+ if (tkind_zero != tkind)
+ {
+ if (integer_zerop (s))
+ tkind = tkind_zero;
+ else if (integer_nonzerop (s))
+ tkind_zero = tkind;
+ }
break;
case OMP_CLAUSE_TO:
tkind = GOMP_MAP_TO;
+ tkind_zero = tkind;
break;
case OMP_CLAUSE_FROM:
tkind = GOMP_MAP_FROM;
+ tkind_zero = tkind;
break;
default:
gcc_unreachable ();
}
gcc_checking_assert (tkind
< (HOST_WIDE_INT_C (1U) << talign_shift));
+ gcc_checking_assert (tkind_zero
+ < (HOST_WIDE_INT_C (1U) << talign_shift));
talign = ceil_log2 (talign);
tkind |= talign << talign_shift;
+ tkind_zero |= talign << talign_shift;
gcc_checking_assert (tkind
<= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
- CONSTRUCTOR_APPEND_ELT (vkind, purpose,
- build_int_cstu (tkind_type, tkind));
+ gcc_checking_assert (tkind_zero
+ <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+ if (tkind == tkind_zero)
+ x = build_int_cstu (tkind_type, tkind);
+ else
+ {
+ TREE_STATIC (TREE_VEC_ELT (t, 2)) = 0;
+ x = build3 (COND_EXPR, tkind_type,
+ fold_build2 (EQ_EXPR, boolean_type_node,
+ unshare_expr (s), size_zero_node),
+ build_int_cstu (tkind_type, tkind_zero),
+ build_int_cstu (tkind_type, tkind));
+ }
+ CONSTRUCTOR_APPEND_ELT (vkind, purpose, x);
if (nc && nc != c)
c = nc;
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
if (is_reference (ovar))
talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
@@ -13157,7 +13362,24 @@ lower_omp_target (gimple_stmt_iterator *
talign = DECL_ALIGN_UNIT (ovar);
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
- if (is_reference (var))
+ tkind = GOMP_MAP_FIRSTPRIVATE;
+ type = TREE_TYPE (ovar);
+ if (is_reference (ovar))
+ type = TREE_TYPE (type);
+ if ((INTEGRAL_TYPE_P (type)
+ && TYPE_PRECISION (type) <= POINTER_SIZE)
+ || TREE_CODE (type) == POINTER_TYPE)
+ {
+ tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ tree t = var;
+ if (is_reference (var))
+ t = build_simple_mem_ref (var);
+ if (TREE_CODE (type) != POINTER_TYPE)
+ t = fold_convert (pointer_sized_int_node, t);
+ t = fold_convert (TREE_TYPE (x), t);
+ gimplify_assign (x, t, &ilist);
+ }
+ else if (is_reference (var))
gimplify_assign (x, var, &ilist);
else if (is_gimple_reg (var))
{
@@ -13172,7 +13394,9 @@ lower_omp_target (gimple_stmt_iterator *
var = build_fold_addr_expr (var);
gimplify_assign (x, var, &ilist);
}
- if (is_reference (var))
+ if (tkind == GOMP_MAP_FIRSTPRIVATE_INT)
+ s = size_int (0);
+ else if (is_reference (var))
s = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
else
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
@@ -13182,7 +13406,6 @@ lower_omp_target (gimple_stmt_iterator *
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
- tkind = GOMP_MAP_FIRSTPRIVATE;
gcc_checking_assert (tkind
< (HOST_WIDE_INT_C (1U) << talign_shift));
talign = ceil_log2 (talign);
@@ -13192,6 +13415,40 @@ lower_omp_target (gimple_stmt_iterator *
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
build_int_cstu (tkind_type, tkind));
break;
+
+ case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_IS_DEVICE_PTR:
+ ovar = OMP_CLAUSE_DECL (c);
+ var = lookup_decl_in_outer_ctx (ovar, ctx);
+ x = build_sender_ref (ovar, ctx);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+ tkind = GOMP_MAP_USE_DEVICE_PTR;
+ else
+ tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ type = TREE_TYPE (ovar);
+ if (TREE_CODE (type) == ARRAY_TYPE)
+ var = build_fold_addr_expr (var);
+ else
+ {
+ if (is_reference (ovar))
+ {
+ type = TREE_TYPE (type);
+ if (TREE_CODE (type) != ARRAY_TYPE)
+ var = build_simple_mem_ref (var);
+ var = fold_convert (TREE_TYPE (x), var);
+ }
+ }
+ gimplify_assign (x, var, &ilist);
+ s = size_int (0);
+ purpose = size_int (map_idx++);
+ CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
+ gcc_checking_assert (tkind
+ < (HOST_WIDE_INT_C (1U) << talign_shift));
+ gcc_checking_assert (tkind
+ <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+ CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+ build_int_cstu (tkind_type, tkind));
+ break;
}
gcc_assert (map_idx == map_cnt);
@@ -13200,21 +13457,22 @@ lower_omp_target (gimple_stmt_iterator *
= 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 (TREE_TYPE (TREE_VEC_ELT (t, 1)),
- NULL);
- TREE_THIS_VOLATILE (clobber) = 1;
- gimple_seq_add_stmt (&olist,
- gimple_build_assign (TREE_VEC_ELT (t, 1),
- clobber));
- }
+ for (int i = 1; i <= 2; i++)
+ if (!TREE_STATIC (TREE_VEC_ELT (t, i)))
+ {
+ gimple_seq initlist = NULL;
+ force_gimple_operand (build1 (DECL_EXPR, void_type_node,
+ TREE_VEC_ELT (t, i)),
+ &initlist, true, NULL_TREE);
+ gimple_seq_add_seq (&ilist, initlist);
+
+ tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, i)),
+ NULL);
+ TREE_THIS_VOLATILE (clobber) = 1;
+ gimple_seq_add_stmt (&olist,
+ gimple_build_assign (TREE_VEC_ELT (t, i),
+ clobber));
+ }
tree clobber = build_constructor (ctx->record_type, NULL);
TREE_THIS_VOLATILE (clobber) = 1;
@@ -13237,22 +13495,64 @@ lower_omp_target (gimple_stmt_iterator *
gimple_build_assign (ctx->receiver_decl, t));
}
- if (offloaded)
+ if (offloaded || data_region)
{
+ tree prev = NULL_TREE;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
- tree var;
+ tree var, x;
default:
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
if (is_reference (var)
|| is_gimple_reg_type (TREE_TYPE (var)))
{
tree new_var = lookup_decl (var, ctx);
- tree x = build_receiver_ref (var, !is_reference (var), ctx);
+ tree type;
+ type = TREE_TYPE (var);
+ if (is_reference (var))
+ type = TREE_TYPE (type);
+ if ((INTEGRAL_TYPE_P (type)
+ && TYPE_PRECISION (type) <= POINTER_SIZE)
+ || TREE_CODE (type) == POINTER_TYPE)
+ {
+ x = build_receiver_ref (var, false, ctx);
+ if (TREE_CODE (type) != POINTER_TYPE)
+ x = fold_convert (pointer_sized_int_node, x);
+ x = fold_convert (type, x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+ fb_rvalue);
+ if (is_reference (var))
+ {
+ tree v = create_tmp_var_raw (type, get_name (var));
+ gimple_add_tmp_var (v);
+ TREE_ADDRESSABLE (v) = 1;
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (v, x));
+ x = build_fold_addr_expr (v);
+ }
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
+ else
+ {
+ x = build_receiver_ref (var, !is_reference (var), ctx);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+ fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
+ }
+ else if (is_variable_sized (var))
+ {
+ tree pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ gcc_assert (DECL_P (pvar));
+ tree new_var = lookup_decl (pvar, ctx);
+ x = build_receiver_ref (var, false, ctx);
gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
gimple_seq_add_stmt (&new_body,
gimple_build_assign (new_var, x));
@@ -13264,23 +13564,22 @@ lower_omp_target (gimple_stmt_iterator *
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
tree new_var = lookup_decl (var, ctx);
- tree x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
+ x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
if (TREE_CONSTANT (x))
{
- const char *name = NULL;
- if (DECL_NAME (var))
- name = IDENTIFIER_POINTER (DECL_NAME (new_var));
-
x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)),
- name);
+ get_name (var));
gimple_add_tmp_var (x);
TREE_ADDRESSABLE (x) = 1;
x = build_fold_addr_expr_loc (clause_loc, x);
}
else
{
- tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA);
- x = build_call_expr_loc (clause_loc, atmp, 1, x);
+ tree atmp
+ = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+ tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+ tree al = size_int (TYPE_ALIGN (rtype));
+ x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
}
x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
@@ -13289,9 +13588,169 @@ lower_omp_target (gimple_stmt_iterator *
gimple_build_assign (new_var, x));
}
break;
+ case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_IS_DEVICE_PTR:
+ var = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+ x = build_sender_ref (var, ctx);
+ else
+ x = build_receiver_ref (var, false, ctx);
+ if (is_variable_sized (var))
+ {
+ tree pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ gcc_assert (DECL_P (pvar));
+ tree new_var = lookup_decl (pvar, ctx);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
+ else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ {
+ tree new_var = lookup_decl (var, ctx);
+ new_var = DECL_VALUE_EXPR (new_var);
+ gcc_assert (TREE_CODE (new_var) == MEM_REF);
+ new_var = TREE_OPERAND (new_var, 0);
+ gcc_assert (DECL_P (new_var));
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
+ else
+ {
+ tree type = TREE_TYPE (var);
+ tree new_var = lookup_decl (var, ctx);
+ if (is_reference (var))
+ {
+ type = TREE_TYPE (type);
+ if (TREE_CODE (type) != ARRAY_TYPE)
+ {
+ tree v = create_tmp_var_raw (type, get_name (var));
+ gimple_add_tmp_var (v);
+ TREE_ADDRESSABLE (v) = 1;
+ x = fold_convert (type, x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val,
+ fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (v, x));
+ x = build_fold_addr_expr (v);
+ }
+ }
+ x = fold_convert (TREE_TYPE (new_var), x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
+ break;
+ }
+ /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass,
+ so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
+ are already handled. */
+ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ tree var;
+ default:
+ break;
+ case OMP_CLAUSE_MAP:
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ gcc_assert (prev);
+ var = OMP_CLAUSE_DECL (c);
+ if (DECL_SIZE (var)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ {
+ tree var2 = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
+ var2 = TREE_OPERAND (var2, 0);
+ gcc_assert (DECL_P (var2));
+ var = var2;
+ }
+ tree new_var = lookup_decl (var, ctx), x;
+ tree type = TREE_TYPE (new_var);
+ bool is_ref = is_reference (var);
+ bool ref_to_array = false;
+ if (is_ref)
+ {
+ type = TREE_TYPE (type);
+ if (TREE_CODE (type) == ARRAY_TYPE)
+ {
+ type = build_pointer_type (type);
+ ref_to_array = true;
+ }
+ }
+ else if (TREE_CODE (type) == ARRAY_TYPE)
+ {
+ tree decl2 = DECL_VALUE_EXPR (new_var);
+ gcc_assert (TREE_CODE (decl2) == MEM_REF);
+ decl2 = TREE_OPERAND (decl2, 0);
+ gcc_assert (DECL_P (decl2));
+ new_var = decl2;
+ type = TREE_TYPE (new_var);
+ }
+ x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx);
+ x = fold_convert_loc (clause_loc, type, x);
+ if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
+ {
+ tree bias = OMP_CLAUSE_SIZE (c);
+ if (DECL_P (bias))
+ bias = lookup_decl (bias, ctx);
+ bias = fold_convert_loc (clause_loc, sizetype, bias);
+ bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype,
+ bias);
+ x = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR,
+ TREE_TYPE (x), x, bias);
+ }
+ if (ref_to_array)
+ x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ if (is_ref && !ref_to_array)
+ {
+ tree t = create_tmp_var_raw (type, get_name (var));
+ gimple_add_tmp_var (t);
+ TREE_ADDRESSABLE (t) = 1;
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (t, x));
+ x = build_fold_addr_expr_loc (clause_loc, t);
+ }
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ prev = NULL_TREE;
+ }
+ else if (OMP_CLAUSE_CHAIN (c)
+ && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+ == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ prev = c;
+ break;
+ case OMP_CLAUSE_PRIVATE:
+ var = OMP_CLAUSE_DECL (c);
+ if (is_variable_sized (var))
+ {
+ location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ tree new_var = lookup_decl (var, ctx);
+ tree pvar = DECL_VALUE_EXPR (var);
+ gcc_assert (TREE_CODE (pvar) == INDIRECT_REF);
+ pvar = TREE_OPERAND (pvar, 0);
+ gcc_assert (DECL_P (pvar));
+ tree new_pvar = lookup_decl (pvar, ctx);
+ tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+ tree al = size_int (DECL_ALIGN (var));
+ tree x = TYPE_SIZE_UNIT (TREE_TYPE (new_var));
+ x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
+ x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_pvar, x));
+ }
+ break;
}
gimple_seq_add_seq (&new_body, tgt_body);
- new_body = maybe_catch_exception (new_body);
+ if (offloaded)
+ new_body = maybe_catch_exception (new_body);
}
else if (data_region)
new_body = tgt_body;
@@ -639,6 +639,9 @@ dump_omp_clause (pretty_printer *pp, tre
case GOMP_MAP_RELEASE:
pp_string (pp, "release");
break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ pp_string (pp, "firstprivate");
+ break;
default:
gcc_unreachable ();
}
@@ -649,7 +652,9 @@ dump_omp_clause (pretty_printer *pp, tre
if (OMP_CLAUSE_SIZE (clause))
{
if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (clause)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER))
pp_string (pp, " [pointer assign, bias: ");
else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET)
@@ -649,7 +649,7 @@ extern tree c_begin_omp_task (void);
extern tree c_finish_omp_task (location_t, tree, tree);
extern void c_finish_omp_cancel (location_t, tree);
extern void c_finish_omp_cancellation_point (location_t, tree);
-extern tree c_finish_omp_clauses (tree, bool = false);
+extern tree c_finish_omp_clauses (tree, bool, bool = false);
extern tree c_build_va_arg (location_t, tree, tree);
extern tree c_finish_transaction (location_t, tree, int);
extern bool c_tree_equal (tree, tree);
@@ -12435,7 +12435,7 @@ c_parser_oacc_all_clauses (c_parser *par
c_parser_skip_to_pragma_eol (parser);
if (finish_p)
- return c_finish_omp_clauses (clauses);
+ return c_finish_omp_clauses (clauses, false);
return clauses;
}
@@ -12720,8 +12720,8 @@ c_parser_omp_all_clauses (c_parser *pars
if (finish_p)
{
if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0)
- return c_finish_omp_clauses (clauses, true);
- return c_finish_omp_clauses (clauses);
+ return c_finish_omp_clauses (clauses, true, true);
+ return c_finish_omp_clauses (clauses, true);
}
return clauses;
@@ -12755,7 +12755,7 @@ c_parser_oacc_cache (location_t loc, c_p
tree stmt, clauses;
clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
- clauses = c_finish_omp_clauses (clauses);
+ clauses = c_finish_omp_clauses (clauses, false);
c_parser_skip_to_pragma_eol (parser);
@@ -13902,7 +13902,7 @@ omp_split_clauses (location_t loc, enum
c_omp_split_clauses (loc, code, mask, clauses, cclauses);
for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
if (cclauses[i])
- cclauses[i] = c_finish_omp_clauses (cclauses[i]);
+ cclauses[i] = c_finish_omp_clauses (cclauses[i], true);
}
/* OpenMP 4.0:
@@ -14668,9 +14668,10 @@ c_parser_omp_target_data (location_t loc
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
- case GOMP_MAP_POINTER:
map_seen = 3;
break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -14800,9 +14801,10 @@ c_parser_omp_target_enter_data (location
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALLOC:
- case GOMP_MAP_POINTER:
map_seen = 3;
break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -14885,9 +14887,10 @@ c_parser_omp_target_exit_data (location_
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
- case GOMP_MAP_POINTER:
map_seen = 3;
break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -15016,6 +15019,7 @@ c_parser_omp_target (c_parser *parser, e
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
OMP_TARGET_BODY (stmt) = block;
+ OMP_TARGET_COMBINED (stmt) = 1;
add_stmt (stmt);
pc = &OMP_TARGET_CLAUSES (stmt);
goto check_clauses;
@@ -15078,7 +15082,7 @@ check_clauses:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
- case GOMP_MAP_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -16379,7 +16383,7 @@ c_parser_cilk_for (c_parser *parser, tre
tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE);
OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR;
OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain;
- clauses = c_finish_omp_clauses (clauses);
+ clauses = c_finish_omp_clauses (clauses, false);
tree block = c_begin_compound_stmt (true);
tree sb = push_stmt_list ();
@@ -16444,7 +16448,7 @@ c_parser_cilk_for (c_parser *parser, tre
OMP_CLAUSE_OPERAND (c, 0)
= cilk_for_number_of_iterations (omp_for);
OMP_CLAUSE_CHAIN (c) = clauses;
- OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c);
+ OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true);
add_stmt (omp_par);
}
@@ -11850,7 +11850,7 @@ handle_omp_array_sections_1 (tree c, tre
/* Handle array sections for clause C. */
static bool
-handle_omp_array_sections (tree c)
+handle_omp_array_sections (tree c, bool is_omp)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
@@ -12030,9 +12030,26 @@ handle_omp_array_sections (tree c)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
+ if (is_omp)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+ break;
+ default:
+ break;
+ }
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
- if (!c_mark_addressable (t))
+ OMP_CLAUSE_SET_MAP_KIND (c2, is_omp
+ ? GOMP_MAP_FIRSTPRIVATE_POINTER
+ : GOMP_MAP_POINTER);
+ if (!is_omp && !c_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
@@ -12097,7 +12114,7 @@ c_find_omp_placeholder_r (tree *tp, int
Remove any elements from the list that are invalid. */
tree
-c_finish_omp_clauses (tree clauses, bool declare_simd)
+c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
bitmap_head aligned_head, map_head;
@@ -12136,7 +12153,7 @@ c_finish_omp_clauses (tree clauses, bool
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c))
+ if (handle_omp_array_sections (c, is_omp))
{
remove = true;
break;
@@ -12496,7 +12513,7 @@ c_finish_omp_clauses (tree clauses, bool
}
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c))
+ if (handle_omp_array_sections (c, is_omp))
remove = true;
break;
}
@@ -12519,7 +12536,7 @@ c_finish_omp_clauses (tree clauses, bool
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c))
+ if (handle_omp_array_sections (c, is_omp))
remove = true;
else
{
@@ -12556,6 +12573,8 @@ c_finish_omp_clauses (tree clauses, bool
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ || (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FORCE_DEVICEPTR)))
&& !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
{
@@ -12624,10 +12643,11 @@ c_finish_omp_clauses (tree clauses, bool
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
+ if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE
+ && TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c),
- "%qs variable is not a pointer",
+ "%qs variable is neither a pointer nor an array",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
@@ -32276,27 +32276,28 @@ cp_parser_omp_target_data (cp_parser *pa
for (tree *pc = &clauses; *pc;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
- switch (OMP_CLAUSE_MAP_KIND (*pc))
- {
- case GOMP_MAP_TO:
- case GOMP_MAP_ALWAYS_TO:
- case GOMP_MAP_FROM:
- case GOMP_MAP_ALWAYS_FROM:
- case GOMP_MAP_TOFROM:
- case GOMP_MAP_ALWAYS_TOFROM:
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_POINTER:
- map_seen = 3;
- break;
- default:
- map_seen |= 1;
- error_at (OMP_CLAUSE_LOCATION (*pc),
- "%<#pragma omp target data%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
- *pc = OMP_CLAUSE_CHAIN (*pc);
- continue;
- }
+ switch (OMP_CLAUSE_MAP_KIND (*pc))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_ALLOC:
+ map_seen = 3;
+ break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ break;
+ default:
+ map_seen |= 1;
+ error_at (OMP_CLAUSE_LOCATION (*pc),
+ "%<#pragma omp target data%> with map-type other "
+ "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+ "on %<map%> clause");
+ *pc = OMP_CLAUSE_CHAIN (*pc);
+ continue;
+ }
pc = &OMP_CLAUSE_CHAIN (*pc);
}
@@ -32370,22 +32371,23 @@ cp_parser_omp_target_enter_data (cp_pars
for (tree *pc = &clauses; *pc;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
- switch (OMP_CLAUSE_MAP_KIND (*pc))
- {
- case GOMP_MAP_TO:
- case GOMP_MAP_ALWAYS_TO:
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_POINTER:
- map_seen = 3;
- break;
- default:
- map_seen |= 1;
- error_at (OMP_CLAUSE_LOCATION (*pc),
- "%<#pragma omp target enter data%> with map-type other "
- "than %<to%> or %<alloc%> on %<map%> clause");
- *pc = OMP_CLAUSE_CHAIN (*pc);
- continue;
- }
+ switch (OMP_CLAUSE_MAP_KIND (*pc))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALLOC:
+ map_seen = 3;
+ break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ break;
+ default:
+ map_seen |= 1;
+ error_at (OMP_CLAUSE_LOCATION (*pc),
+ "%<#pragma omp target enter data%> with map-type other "
+ "than %<to%> or %<alloc%> on %<map%> clause");
+ *pc = OMP_CLAUSE_CHAIN (*pc);
+ continue;
+ }
pc = &OMP_CLAUSE_CHAIN (*pc);
}
@@ -32455,24 +32457,25 @@ cp_parser_omp_target_exit_data (cp_parse
for (tree *pc = &clauses; *pc;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
- switch (OMP_CLAUSE_MAP_KIND (*pc))
- {
- case GOMP_MAP_FROM:
- case GOMP_MAP_ALWAYS_FROM:
- case GOMP_MAP_RELEASE:
- case GOMP_MAP_DELETE:
- case GOMP_MAP_POINTER:
- map_seen = 3;
- break;
- default:
- map_seen |= 1;
- error_at (OMP_CLAUSE_LOCATION (*pc),
- "%<#pragma omp target exit data%> with map-type other "
- "than %<from%>, %<release%> or %<delete%> on %<map%>"
- " clause");
- *pc = OMP_CLAUSE_CHAIN (*pc);
- continue;
- }
+ switch (OMP_CLAUSE_MAP_KIND (*pc))
+ {
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ map_seen = 3;
+ break;
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ break;
+ default:
+ map_seen |= 1;
+ error_at (OMP_CLAUSE_LOCATION (*pc),
+ "%<#pragma omp target exit data%> with map-type other "
+ "than %<from%>, %<release%> or %<delete%> on %<map%>"
+ " clause");
+ *pc = OMP_CLAUSE_CHAIN (*pc);
+ continue;
+ }
pc = &OMP_CLAUSE_CHAIN (*pc);
}
@@ -32637,6 +32640,7 @@ cp_parser_omp_target (cp_parser *parser,
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
OMP_TARGET_BODY (stmt) = body;
+ OMP_TARGET_COMBINED (stmt) = 1;
add_stmt (stmt);
pc = &OMP_TARGET_CLAUSES (stmt);
goto check_clauses;
@@ -32697,7 +32701,7 @@ check_clauses:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
- case GOMP_MAP_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -4650,7 +4650,7 @@ handle_omp_array_sections_1 (tree c, tre
/* Handle array sections for clause C. */
static bool
-handle_omp_array_sections (tree c)
+handle_omp_array_sections (tree c, bool is_omp)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
@@ -4826,10 +4826,26 @@ handle_omp_array_sections (tree c)
OMP_CLAUSE_SIZE (c) = size;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
+ if (is_omp)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+ break;
+ default:
+ break;
+ }
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
- if (!cxx_mark_addressable (t))
+ OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER
+ : GOMP_MAP_POINTER);
+ if (!is_omp && !cxx_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
@@ -4847,7 +4863,8 @@ handle_omp_array_sections (tree c)
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
ptr = OMP_CLAUSE_DECL (c2);
- if (TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
+ if (!is_omp
+ && TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
&& POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
{
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
@@ -5569,7 +5586,7 @@ finish_omp_clauses (tree clauses, bool a
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c))
+ if (handle_omp_array_sections (c, allow_fields))
{
remove = true;
break;
@@ -6155,7 +6172,7 @@ finish_omp_clauses (tree clauses, bool a
}
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c))
+ if (handle_omp_array_sections (c, allow_fields))
remove = true;
break;
}
@@ -6189,7 +6206,7 @@ finish_omp_clauses (tree clauses, bool a
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c))
+ if (handle_omp_array_sections (c, allow_fields))
remove = true;
else
{
@@ -6242,7 +6259,9 @@ finish_omp_clauses (tree clauses, bool a
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER)))
&& !type_dependent_expression_p (t)
&& !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
== REFERENCE_TYPE)
@@ -6428,12 +6447,14 @@ finish_omp_clauses (tree clauses, bool a
{
tree type = TREE_TYPE (t);
if (TREE_CODE (type) != POINTER_TYPE
+ && TREE_CODE (type) != ARRAY_TYPE
&& (TREE_CODE (type) != REFERENCE_TYPE
- || TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE))
+ || (TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE
+ && TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE)))
{
error_at (OMP_CLAUSE_LOCATION (c),
- "%qs variable is not a pointer or reference "
- "to pointer",
+ "%qs variable is neither a pointer, nor an array"
+ "nor reference to pointer or array",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
@@ -74,6 +74,17 @@ enum gomp_map_kind
GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
+ /* Similarly, but store the value in the pointer rather than
+ pointed by the pointer. */
+ GOMP_MAP_FIRSTPRIVATE_INT = (GOMP_MAP_FLAG_SPECIAL | 1),
+ /* Pointer translate host address into device address and copy that
+ back to host. */
+ GOMP_MAP_USE_DEVICE_PTR = (GOMP_MAP_FLAG_SPECIAL | 2),
+ /* Allocate a zero length array section. Prefer next non-zero length
+ mapping over previous non-zero length mapping over zero length mapping
+ at the address. If not already mapped, do nothing (and pointer translate
+ to NULL). */
+ GOMP_MAP_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL | 3),
/* Allocate. */
GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC),
/* ..., and copy to device. */
@@ -95,7 +106,11 @@ enum gomp_map_kind
GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC,
/* Decrement usage count and deallocate if zero. */
GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_ALWAYS
- | GOMP_MAP_FORCE_DEALLOC)
+ | GOMP_MAP_FORCE_DEALLOC),
+
+ /* Internal to GCC, not used in libgomp. */
+ /* Do not map, but pointer assign a pointer instead. */
+ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1)
};
#define GOMP_MAP_COPY_TO_P(X) \
@@ -647,11 +647,9 @@ struct target_var_desc {
bool copy_from;
/* True if data always should be copied from device to host at the end. */
bool always_copy_from;
- /* Used for unmapping of array sections, can be nonzero only when
- always_copy_from is true. */
+ /* Relative offset against key host_start. */
uintptr_t offset;
- /* Used for unmapping of array sections, can be less than the size of the
- whole object only when always_copy_from is true. */
+ /* Actual length. */
uintptr_t length;
};
@@ -142,7 +142,26 @@ resolve_device (int device_id)
}
-/* Handle the case where splay_tree_lookup found oldn for newn.
+static inline splay_tree_key
+gomp_map_lookup (splay_tree mem_map, splay_tree_key key)
+{
+ if (key->host_start != key->host_end)
+ return splay_tree_lookup (mem_map, key);
+
+ key->host_end++;
+ splay_tree_key n = splay_tree_lookup (mem_map, key);
+ key->host_end--;
+ if (n)
+ return n;
+ key->host_start--;
+ n = splay_tree_lookup (mem_map, key);
+ key->host_start++;
+ if (n)
+ return n;
+ return splay_tree_lookup (mem_map, key);
+}
+
+/* Handle the case where gomp_map_lookup found oldn for newn.
Helper function of gomp_map_vars. */
static inline void
@@ -204,20 +223,8 @@ gomp_map_pointer (struct target_mem_desc
}
/* Add bias to the pointer value. */
cur_node.host_start += bias;
- cur_node.host_end = cur_node.host_start + 1;
- splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
- if (n == NULL)
- {
- /* Could be possibly zero size array section. */
- cur_node.host_end--;
- n = splay_tree_lookup (mem_map, &cur_node);
- if (n == NULL)
- {
- cur_node.host_start--;
- n = splay_tree_lookup (mem_map, &cur_node);
- cur_node.host_start++;
- }
- }
+ cur_node.host_end = cur_node.host_start;
+ splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
if (n == NULL)
{
gomp_mutex_unlock (&devicep->lock);
@@ -271,9 +278,29 @@ gomp_map_vars (struct gomp_device_descr
for (i = 0; i < mapnum; i++)
{
int kind = get_kind (short_mapkind, kinds, i);
- if (hostaddrs[i] == NULL)
+ if (hostaddrs[i] == NULL
+ || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT)
{
tgt->list[i].key = NULL;
+ tgt->list[i].offset = ~(uintptr_t) 0;
+ continue;
+ }
+ else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start;
+ splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("use_device_ptr pointer wasn't mapped");
+ }
+ cur_node.host_start -= n->host_start;
+ hostaddrs[i]
+ = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start);
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = ~(uintptr_t) 0;
continue;
}
cur_node.host_start = (uintptr_t) hostaddrs[i];
@@ -293,7 +320,19 @@ gomp_map_vars (struct gomp_device_descr
has_firstprivate = true;
continue;
}
- splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ splay_tree_key n;
+ if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
+ {
+ n = gomp_map_lookup (mem_map, &cur_node);
+ if (!n)
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = ~(uintptr_t) 1;
+ continue;
+ }
+ }
+ else
+ n = splay_tree_lookup (mem_map, &cur_node);
if (n)
gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
kind & typemask);
@@ -386,6 +425,15 @@ gomp_map_vars (struct gomp_device_descr
tgt_size += len;
continue;
}
+ switch (kind & typemask)
+ {
+ case GOMP_MAP_FIRSTPRIVATE_INT:
+ case GOMP_MAP_USE_DEVICE_PTR:
+ case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+ continue;
+ default:
+ break;
+ }
splay_tree_key k = &array->key;
k->host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
@@ -518,15 +566,18 @@ gomp_map_vars (struct gomp_device_descr
{
if (tgt->list[i].key == NULL)
{
- if (hostaddrs[i] == NULL)
- cur_node.tgt_offset = (uintptr_t) NULL;
+ if (tgt->list[i].offset == ~(uintptr_t) 0)
+ cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
+ else if (tgt->list[i].offset == ~(uintptr_t) 1)
+ cur_node.tgt_offset = 0;
else
cur_node.tgt_offset = tgt->tgt_start
+ tgt->list[i].offset;
}
else
cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset;
+ + tgt->list[i].key->tgt_offset
+ + tgt->list[i].offset;
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -1052,7 +1103,38 @@ GOMP_target_41 (int device, void (*fn) (
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
- return gomp_target_fallback (fn, hostaddrs);
+ {
+ size_t i, tgt_align = 0, tgt_size = 0;
+ char *tgt = NULL;
+ for (i = 0; i < mapnum; i++)
+ if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+ {
+ size_t align = (size_t) 1 << (kinds[i] >> 8);
+ if (tgt_align < align)
+ tgt_align = align;
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += sizes[i];
+ }
+ if (tgt_align)
+ {
+ tgt = gomp_alloca (tgt_size + tgt_align - 1);
+ uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
+ if (al)
+ tgt += tgt_align - al;
+ tgt_size = 0;
+ for (i = 0; i < mapnum; i++)
+ if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+ {
+ size_t align = (size_t) 1 << (kinds[i] >> 8);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
+ hostaddrs[i] = tgt + tgt_size;
+ tgt_size = tgt_size + sizes[i];
+ }
+ }
+ gomp_target_fallback (fn, hostaddrs);
+ return;
+ }
void *fn_addr = gomp_get_target_fn_addr (devicep, fn);
@@ -1289,20 +1371,8 @@ omp_target_is_present (void *ptr, size_t
struct splay_tree_key_s cur_node;
cur_node.host_start = (uintptr_t) ptr + offset;
- cur_node.host_end = cur_node.host_start + 1;
- splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
- if (n == NULL)
- {
- /* Could be possibly zero size array section. */
- cur_node.host_end--;
- n = splay_tree_lookup (mem_map, &cur_node);
- if (n == NULL)
- {
- cur_node.host_start--;
- n = splay_tree_lookup (mem_map, &cur_node);
- cur_node.host_start++;
- }
- }
+ cur_node.host_end = cur_node.host_start;
+ splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
int ret = n != NULL;
gomp_mutex_unlock (&devicep->lock);
return ret;
@@ -1524,7 +1594,7 @@ omp_target_associate_ptr (void *host_ptr
cur_node.host_start = (uintptr_t) host_ptr;
cur_node.host_end = cur_node.host_start + size;
- splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
if (n)
{
if (n->tgt->tgt_start + n->tgt_offset
@@ -1584,13 +1654,8 @@ omp_target_disassociate_ptr (void *ptr,
int ret = EINVAL;
cur_node.host_start = (uintptr_t) ptr;
- cur_node.host_end = cur_node.host_start + 1;
- splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
- if (n == NULL)
- {
- cur_node.host_end--;
- n = splay_tree_lookup (mem_map, &cur_node);
- }
+ cur_node.host_end = cur_node.host_start;
+ splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
if (n
&& n->host_start == cur_node.host_start
&& n->refcount == REFCOUNT_INFINITY
@@ -33,7 +33,8 @@ fn2 (int x, double (&dr) [1024], double
int j;
fn1 (hr + 2 * x, ir + 2 * x, x);
#pragma omp target map(to: br[:x], cr[0:x], dr[x:x], er[x:x]) \
- map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x])
+ map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x]) \
+ map(tofrom: s)
#pragma omp parallel for reduction(+:s)
for (j = 0; j < x; j++)
s += br[j] * cr[j] + dr[x + j] + er[x + j]
@@ -0,0 +1,90 @@
+extern "C" void abort ();
+
+void
+foo (int *x, int *&y, int (&z)[15])
+{
+ int a[10], b[15], err, i;
+ for (i = 0; i < 10; i++)
+ a[i] = 7 * i;
+ for (i = 0; i < 15; i++)
+ b[i] = 8 * i;
+ #pragma omp target map(to:x[5:10], y[5:10], z[5:10], a[0:10], b[5:10]) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 10; i++)
+ if (x[5 + i] != 20 + 4 * i
+ || y[5 + i] != 25 + 5 * i
+ || z[5 + i] != 30 + 6 * i
+ || a[i] != 7 * i
+ || b[5 + i] != 40 + 8 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+void
+bar (int n, int v)
+{
+ int a[n], b[n], c[n], d[n], e[n], err, i;
+ int (*x)[n] = &c;
+ int (*y2)[n] = &d;
+ int (*&y)[n] = y2;
+ int (&z)[n] = e;
+ for (i = 0; i < n; i++)
+ {
+ (*x)[i] = 4 * i;
+ (*y)[i] = 5 * i;
+ z[i] = 6 * i;
+ a[i] = 7 * i;
+ b[i] = 8 * i;
+ }
+ #pragma omp target map(to:x[0][5:10], y[0][5:10], z[5:10], a[0:10], b[5:10]) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 10; i++)
+ if ((*x)[5 + i] != 20 + 4 * i
+ || (*y)[5 + i] != 25 + 5 * i
+ || z[5 + i] != 30 + 6 * i
+ || a[i] != 7 * i
+ || b[5 + i] != 40 + 8 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ {
+ (*x)[i] = 9 * i;
+ (*y)[i] = 10 * i;
+ z[i] = 11 * i;
+ a[i] = 12 * i;
+ b[i] = 13 * i;
+ }
+ #pragma omp target map(to:x[0][v:v+5], y[0][v:v+5], z[v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 10; i++)
+ if ((*x)[5 + i] != 45 + 9 * i
+ || (*y)[5 + i] != 50 + 10 * i
+ || z[5 + i] != 55 + 11 * i
+ || a[i] != 12 * i
+ || b[5 + i] != 65 + 13 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ int x[15], y2[15], z[15], *y = y2, i;
+ for (i = 0; i < 15; i++)
+ {
+ x[i] = 4 * i;
+ y[i] = 5 * i;
+ z[i] = 6 * i;
+ }
+ foo (x, y, z);
+ bar (15, 5);
+}
@@ -0,0 +1,58 @@
+extern "C" void abort ();
+struct S { int a; };
+#ifdef __SIZEOF_INT128__
+typedef __int128 T;
+#else
+typedef long long int T;
+#endif
+
+void
+foo (T a, int b, struct S c)
+{
+ int err;
+ #pragma omp target firstprivate (a, b, c) map(from:err)
+ {
+ err = 0;
+ if (a != 131 || b != 276 || c.a != 59)
+ err = 1;
+ a = 936;
+ b = 27;
+ c.a = 98;
+ if (a != 936 || b != 27 || c.a != 98)
+ err = 1;
+ }
+ if (err || a != 131 || b != 276 || c.a != 59)
+ abort ();
+}
+
+void
+bar (T &a, int &b, struct S &c)
+{
+ int err;
+ #pragma omp target firstprivate (a, b, c) map(from:err)
+ {
+ err = 0;
+ if (a != 131 || b != 276 || c.a != 59)
+ err = 1;
+ a = 936;
+ b = 27;
+ c.a = 98;
+ if (a != 936 || b != 27 || c.a != 98)
+ err = 1;
+ }
+ if (err || a != 131 || b != 276 || c.a != 59)
+ abort ();
+}
+
+int
+main ()
+{
+ T a = 131;
+ int b = 276;
+ struct S c;
+ c.a = 59;
+ foo (a, b, c);
+ bar (a, b, c);
+ if (a != 131 || b != 276 || c.a != 59)
+ abort ();
+}
@@ -0,0 +1,73 @@
+extern "C" void abort (void);
+
+void
+foo (int *&p, int (&s)[5], int n)
+{
+ int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
+ int *r = a + 1, *q = p - 1, i, err;
+ for (i = 0; i < n; i++)
+ b[i] = 9 + i;
+ #pragma omp target data map(to:a)
+ #pragma omp target data use_device_ptr(r) map(from:err)
+ #pragma omp target is_device_ptr(r) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (r[i - 1] != 7 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:q[:4])
+ #pragma omp target data use_device_ptr(p) map(from:err)
+ #pragma omp target is_device_ptr(p) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (p[i - 1] != i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:b)
+ #pragma omp target data use_device_ptr(b) map(from:err)
+ #pragma omp target is_device_ptr(b) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (b[i] != 9 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:c)
+ #pragma omp target data use_device_ptr(c) map(from:err)
+ #pragma omp target is_device_ptr(c) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 3; i++)
+ if (c[i] != 20 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:s[:5])
+ #pragma omp target data use_device_ptr(s) map(from:err)
+ #pragma omp target is_device_ptr(s) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 5; i++)
+ if (s[i] != 17 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
+ int *p = a + 1;
+ foo (p, b, 9);
+}
@@ -34,7 +34,7 @@ fn2 (int x, int y, int z)
fn1 (b, c, x);
#pragma omp target data map(to: b)
{
- #pragma omp target map(tofrom: c)
+ #pragma omp target map(tofrom: c, s)
#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)
@@ -52,7 +52,7 @@ 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 target map(to: b, c) map(tofrom:s)
#pragma omp parallel for reduction(+:s)
for (i = 0; i < x; i++)
tgt (), s += b[i] * c[i];
@@ -66,7 +66,8 @@ fn4 (int x, double *p)
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 target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \
+ map(tofrom: s)
#pragma omp parallel for reduction(+:s)
for (i = 0; i < x; i++)
s += b[i] * c[i] + d[x + i] + p[x + i];
@@ -23,7 +23,7 @@ fn2 (int x)
int i;
fn1 (b, c, x);
fn1 (e, d + x, x);
- #pragma omp target map(to: b, c[:x], d[x:x], e)
+ #pragma omp target map(to: b, c[:x], d[x:x], e) map(tofrom: s)
#pragma omp parallel for reduction(+:s)
for (i = 0; i < x; i++)
s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
@@ -38,7 +38,7 @@ fn3 (int x)
int i;
fn1 (b, c, x);
fn1 (e, d, x);
- #pragma omp target
+ #pragma omp target map(tofrom: s)
#pragma omp parallel for reduction(+:s)
for (i = 0; i < x; i++)
s += b[i] * c[i] + d[i];
@@ -56,7 +56,7 @@ fn4 (int x)
#pragma omp target data map(from: b, c[:x], d[x:x], e)
{
#pragma omp target update to(b, c[:x], d[x:x], e)
- #pragma omp target map(c[:x], d[x:x])
+ #pragma omp target map(c[:x], d[x:x], s)
#pragma omp parallel for reduction(+:s)
for (i = 0; i < x; i++)
{
@@ -37,63 +37,63 @@ foo (int f)
abort ();
#pragma omp target data device (d) map (to: h)
{
- #pragma omp target device (d)
+ #pragma omp target device (d) map (h)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5)
abort ();
#pragma omp target update device (d) from (h)
}
#pragma omp target data if (v > 1) map (to: h)
{
- #pragma omp target if (v > 1)
+ #pragma omp target if (v > 1) map(h)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6)
abort ();
#pragma omp target update if (v > 1) from (h)
}
#pragma omp target data device (d) if (v > 1) map (to: h)
{
- #pragma omp target device (d) if (v > 1)
+ #pragma omp target device (d) if (v > 1) map(h)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7)
abort ();
#pragma omp target update device (d) if (v > 1) from (h)
}
#pragma omp target data if (v <= 1) map (to: h)
{
- #pragma omp target if (v <= 1)
+ #pragma omp target if (v <= 1) map (tofrom: h)
if (omp_get_level () != 0 || h++ != 8)
abort ();
#pragma omp target update if (v <= 1) from (h)
}
#pragma omp target data device (d) if (v <= 1) map (to: h)
{
- #pragma omp target device (d) if (v <= 1)
+ #pragma omp target device (d) if (v <= 1) map (h)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9)
abort ();
#pragma omp target update device (d) if (v <= 1) from (h)
}
#pragma omp target data if (0) map (to: h)
{
- #pragma omp target if (0)
+ #pragma omp target if (0) map (h)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10)
abort ();
#pragma omp target update if (0) from (h)
}
#pragma omp target data device (d) if (0) map (to: h)
{
- #pragma omp target device (d) if (0)
+ #pragma omp target device (d) if (0) map (h)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11)
abort ();
#pragma omp target update device (d) if (0) from (h)
}
#pragma omp target data if (1) map (to: h)
{
- #pragma omp target if (1)
+ #pragma omp target if (1) map (tofrom: h)
if (omp_get_level () != 0 || h++ != 12)
abort ();
#pragma omp target update if (1) from (h)
}
#pragma omp target data device (d) if (1) map (to: h)
{
- #pragma omp target device (d) if (1)
+ #pragma omp target device (d) if (1) map (tofrom: h)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13)
abort ();
#pragma omp target update device (d) if (1) from (h)
@@ -0,0 +1,74 @@
+extern void abort (void);
+
+void
+foo (int *x)
+{
+ int a[10], b[15], err, i;
+ for (i = 0; i < 10; i++)
+ a[i] = 7 * i;
+ for (i = 0; i < 15; i++)
+ b[i] = 8 * i;
+ #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 10; i++)
+ if (x[5 + i] != 20 + 4 * i
+ || a[i] != 7 * i
+ || b[5 + i] != 40 + 8 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+void
+bar (int n, int v)
+{
+ int a[n], b[n], c[n], d[n], e[n], err, i;
+ int (*x)[n] = &c;
+ for (i = 0; i < n; i++)
+ {
+ (*x)[i] = 4 * i;
+ a[i] = 7 * i;
+ b[i] = 8 * i;
+ }
+ #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 10; i++)
+ if ((*x)[5 + i] != 20 + 4 * i
+ || a[i] != 7 * i
+ || b[5 + i] != 40 + 8 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ {
+ (*x)[i] = 9 * i;
+ a[i] = 12 * i;
+ b[i] = 13 * i;
+ }
+ #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 10; i++)
+ if ((*x)[5 + i] != 45 + 9 * i
+ || a[i] != 12 * i
+ || b[5 + i] != 65 + 13 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ int x[15], i;
+ for (i = 0; i < 15; i++)
+ x[i] = 4 * i;
+ foo (x);
+ bar (15, 5);
+ return 0;
+}
@@ -0,0 +1,45 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+ int a[n], i, err;
+ for (i = 0; i < n; i++)
+ a[i] = 7 * i;
+ #pragma omp target firstprivate (a) map(from:err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (a[i] != 7 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+void
+bar (int n)
+{
+ int a[n], i, err;
+ #pragma omp target private (a) map(from:err)
+ {
+ #pragma omp parallel for
+ for (i = 0; i < n; i++)
+ a[i] = 7 * i;
+ err = 0;
+ #pragma omp parallel for reduction(|:err)
+ for (i = 0; i < n; i++)
+ if (a[i] != 7 * i)
+ err |= 1;
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ foo (7);
+ bar (7);
+ return 0;
+}
@@ -0,0 +1,99 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+ int a[n], i, err;
+ for (i = 0; i < n; i++)
+ a[i] = 5 * i;
+ #pragma omp target map(to:a) map(from:err) private(i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (a[i] != 5 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ a[i] += i;
+ #pragma omp target map(from:err) private(i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (a[i] != 6 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ a[i] += i;
+ #pragma omp target firstprivate (a) map(from:err) private(i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (a[i] != 7 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ int on = n;
+ #pragma omp target firstprivate (n) map(tofrom: n)
+ {
+ n++;
+ }
+ if (on != n)
+ abort ();
+ #pragma omp target map(tofrom: n) private (n)
+ {
+ n = 25;
+ }
+ if (on != n)
+ abort ();
+ for (i = 0; i < n; i++)
+ a[i] += i;
+ #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (a[i] != 8 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ a[i] += i;
+ #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ if (a[i] != 9 * i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ a[i] += i;
+ #pragma omp target map(tofrom:a) map(from:err) private(a, i)
+ {
+ err = 0;
+ for (i = 0; i < n; i++)
+ a[i] = 7;
+ #pragma omp parallel for reduction(|:err)
+ for (i = 0; i < n; i++)
+ if (a[i] != 7)
+ err |= 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < n; i++)
+ if (a[i] != 10 * i)
+ abort ();
+}
+
+int
+main ()
+{
+ foo (9);
+ return 0;
+}
@@ -0,0 +1,52 @@
+extern void abort (void);
+
+void
+foo (int n)
+{
+ int a[4] = { 0, 1, 2, 3 }, b[n];
+ int *p = a + 1, i, err;
+ for (i = 0; i < n; i++)
+ b[i] = 9 + i;
+ #pragma omp target data map(to:a)
+ #pragma omp target data use_device_ptr(p) map(from:err)
+ #pragma omp target is_device_ptr(p) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (p[i - 1] != i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 4; i++)
+ a[i] = 23 + i;
+ #pragma omp target data map(to:a)
+ #pragma omp target data use_device_ptr(a) map(from:err)
+ #pragma omp target is_device_ptr(a) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (a[i] != 23 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:b)
+ #pragma omp target data use_device_ptr(b) map(from:err)
+ #pragma omp target is_device_ptr(b) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 4; i++)
+ if (b[i] != 9 + i)
+ err = 1;
+ }
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ foo (9);
+ return 0;
+}
@@ -0,0 +1,127 @@
+extern void abort (void);
+
+void
+foo (int *p, int *q, int *r, int n, int m)
+{
+ int i, err, *s = r;
+ #pragma omp target data map(to:p[0:8])
+ {
+ /* For zero length array sections, p points to the start of
+ already mapped range, q to the end of it, and r does not point
+ to an mapped range. */
+ #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1 || q[i - 8] != i + 1)
+ err = 1;
+ if (p + 8 != q || (r != (int *) 0 && r != s))
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Implicit mapping of pointers behaves the same way. */
+ #pragma omp target private(i) map(from:err) firstprivate (s)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1 || q[i - 8] != i + 1)
+ err = 1;
+ if (p + 8 != q || (r != (int *) 0 && r != s))
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* And zero-length array sections, though not known at compile
+ time, behave the same. */
+ #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1 || q[i - 8] != i + 1)
+ err = 1;
+ if (p + 8 != q || (r != (int *) 0 && r != s))
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Non-zero length array sections, though not known at compile,
+ behave differently. */
+ #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r[0] != 10)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:q[0:1])
+ {
+ /* For zero length array sections, p points to the start of
+ already mapped range, q points to the start of another one,
+ and r to the end of the second one. */
+ #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r != q + 1)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Implicit mapping of pointers behaves the same way. */
+ #pragma omp target private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r != q + 1)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* And zero-length array sections, though not known at compile
+ time, behave the same. */
+ #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r != q + 1)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Non-zero length array sections, though not known at compile,
+ behave differently. */
+ #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r[0] != 10)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ }
+ }
+}
+
+int
+main ()
+{
+ int a[32], i;
+ for (i = 0; i < 32; i++)
+ a[i] = i;
+ foo (a + 1, a + 9, a + 10, 0, 1);
+ return 0;
+}
@@ -47,7 +47,7 @@ void gramSchmidt (int Q[][COLS], const i
{
int tmp = 0;
- #pragma omp target
+ #pragma omp target map(tofrom:tmp)
#pragma omp parallel for reduction(+:tmp)
for (i = 0; i < rows; i++)
tmp += (Q[i][k] * Q[i][k]);
@@ -20,7 +20,7 @@ int fib_wrapper (int n)
{
int x = 0;
- #pragma omp target if(n > THRESHOLD)
+ #pragma omp target if(n > THRESHOLD) map(from:x)
x = fib (n);
return x;
@@ -41,7 +41,7 @@ float accum (int k)
int i;
float tmp = 0.0;
- #pragma omp target
+ #pragma omp target map(tofrom:tmp)
#pragma omp parallel for reduction(+:tmp)
for (i = 0; i < N; i++)
tmp += Pfun (i, k);
@@ -48,7 +48,7 @@ float accum ()
int i, k;
float tmp = 0.0;
- #pragma omp target
+ #pragma omp target map(tofrom:tmp)
#pragma omp parallel for reduction(+:tmp)
for (i = 0; i < N; i++)
{
@@ -32,7 +32,7 @@ float dotprod (float B[], float C[], int
int i, i0;
float sum = 0;
- #pragma omp target map(to: B[0:n], C[0:n])
+ #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom: sum)
#pragma omp teams num_teams(num_teams) thread_limit(block_threads) \
reduction(+:sum)
#pragma omp distribute
@@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int
int i;
float sum = 0;
- #pragma omp target teams map(to: B[0:n], C[0:n])
+ #pragma omp target teams map(to: B[0:n], C[0:n]) map(tofrom: sum)
#pragma omp distribute parallel for reduction(+:sum)
for (i = 0; i < n; i++)
sum += B[i] * C[i];
@@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int
int i;
float sum = 0;
- #pragma omp target map(to: B[0:n], C[0:n])
+ #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom:sum)
#pragma omp teams num_teams(8) thread_limit(16)
#pragma omp distribute parallel for reduction(+:sum) \
dist_schedule(static, 1024) \
@@ -10,11 +10,11 @@ int main ()
int b = 0;
int c, d;
- #pragma omp target if(a > 200 && a < 400)
+ #pragma omp target if(a > 200 && a < 400) map(from: c)
c = omp_is_initial_device ();
#pragma omp target data map(to: b) if(a > 200 && a < 400)
- #pragma omp target
+ #pragma omp target map(from: b, d)
{
b = 100;
d = omp_is_initial_device ();
@@ -26,11 +26,11 @@ int main ()
a += 200;
b = 0;
- #pragma omp target if(a > 200 && a < 400)
+ #pragma omp target if(a > 200 && a < 400) map(from: c)
c = omp_is_initial_device ();
#pragma omp target data map(to: b) if(a > 200 && a < 400)
- #pragma omp target
+ #pragma omp target map(from: b, d)
{
b = 100;
d = omp_is_initial_device ();
@@ -42,11 +42,11 @@ int main ()
a += 200;
b = 0;
- #pragma omp target if(a > 200 && a < 400)
+ #pragma omp target if(a > 200 && a < 400) map(from: c)
c = omp_is_initial_device ();
#pragma omp target data map(to: b) if(a > 200 && a < 400)
- #pragma omp target
+ #pragma omp target map(from: b, d)
{
b = 100;
d = omp_is_initial_device ();
@@ -9,7 +9,7 @@ int main ()
int res;
int default_device = omp_get_default_device ();
- #pragma omp target
+ #pragma omp target map(from: res)
res = omp_is_initial_device ();
if (res)
@@ -17,7 +17,7 @@ int main ()
omp_set_default_device (omp_get_num_devices ());
- #pragma omp target
+ #pragma omp target map(from: res)
res = omp_is_initial_device ();
if (!res)