@@ -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,74 @@
+extern void abort ();
+
+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;
+}
@@ -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 gmp_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);
@@ -293,7 +300,7 @@ 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 = gomp_map_lookup (mem_map, &cur_node);
if (n)
gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
kind & typemask);
@@ -392,7 +399,7 @@ gomp_map_vars (struct gomp_device_descr
k->host_end = k->host_start + sizes[i];
else
k->host_end = k->host_start + sizeof (void *);
- splay_tree_key n = splay_tree_lookup (mem_map, k);
+ splay_tree_key n = gomp_map_lookup (mem_map, k);
if (n)
gomp_map_vars_existing (devicep, n, k, &tgt->list[i],
kind & typemask);
@@ -526,7 +533,8 @@ gomp_map_vars (struct gomp_device_descr
}
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
@@ -1289,20 +1297,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 +1520,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 +1580,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
@@ -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;
};
@@ -95,7 +95,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) \
@@ -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);
}
@@ -32697,7 +32700,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;
@@ -4828,8 +4828,9 @@ handle_omp_array_sections (tree c)
return false;
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 +4848,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 +5571,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 +6157,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 +6191,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 +6244,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)
@@ -1445,7 +1445,7 @@ 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
@@ -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)
@@ -156,6 +158,8 @@ 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;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -5821,6 +5825,38 @@ omp_notice_variable (struct gimplify_omp
"a mappable type", decl);
omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
}
+ else if (ctx->target_map_pointers_as_0len_arrays
+ || ctx->target_map_scalars_firstprivate)
+ {
+ bool is_declare_target = 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)
+ omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ 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)))
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_0LEN_ARRAY
+ | flags);
+ else
+ omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ }
else
omp_add_variable (ctx, decl, GOVD_MAP | flags);
}
@@ -6144,6 +6180,13 @@ 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;
+ }
while ((c = *list_p) != NULL)
{
@@ -6319,10 +6362,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 +6389,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 +6708,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 +6827,29 @@ 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_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,
@@ -6915,7 +7006,8 @@ gimplify_adjust_omp_clauses (gimple_seq
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
@@ -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);
@@ -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;
@@ -12031,8 +12031,10 @@ handle_omp_array_sections (tree c)
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
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 +12099,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 +12138,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 +12498,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 +12521,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 +12558,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)))
{
@@ -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),
@@ -15078,7 +15081,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 +16382,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 +16447,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);
}
@@ -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;
@@ -2025,6 +2025,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)
@@ -2201,7 +2216,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)))
{
@@ -12883,6 +12899,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 +12935,25 @@ 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);
+ const char *name = NULL;
+ if (DECL_NAME (new_var))
+ name = IDENTIFIER_POINTER (DECL_NAME (new_var));
+ x = create_tmp_var_raw (type, name);
+ 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 (!maybe_lookup_field (var, ctx))
continue;
@@ -12925,6 +12961,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)
@@ -13044,6 +13081,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)
{
@@ -13239,6 +13280,7 @@ lower_omp_target (gimple_stmt_iterator *
if (offloaded)
{
+ tree prev = NULL_TREE;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
@@ -13290,6 +13332,93 @@ lower_omp_target (gimple_stmt_iterator *
}
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)
+ {
+ const char *name = NULL;
+ if (DECL_NAME (var))
+ name = IDENTIFIER_POINTER (DECL_NAME (new_var));
+
+ tree t = create_tmp_var_raw (type, name);
+ 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;
+ }
gimple_seq_add_seq (&new_body, tgt_body);
new_body = maybe_catch_exception (new_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)