On Mon, Jul 20, 2015 at 08:10:41PM +0200, Jakub Jelinek wrote:
> And here is untested incremental libgomp side of the proposed
> GOMP_MAP_FIRSTPRIVATE_POINTER.
Actually, that seems unnecessary, for the array section maps we already
have there a pointer, so we can easily implement that just on the
compiler side.
Here is a WIP patch.
Unfortunately, in order not to break numerous examples-4/ testcases that
were doing target data map of array sections with target region without
any explicit maps, with the new way where target{, enter, exit} data
no longer map the base pointer, I had to implement the new implicit
pointer mapping semantics (map (alloc:ptr[0:0])) already in this patch.
And, that patch really requires that if there is ptr[0:something] for
something > 0 already mapped that we use the ptr[0:something] mapping rather
than ptr[0:0]. See the libgomp changes for that.
Unfortunately, that occassionally breaks the target8.f90 testcase at -O0,
where we map zero-sized FRAME.6 object which happens to be adjacent to the
array. And that reveals IMNSHO very serious flaw in the current standard
draft, no idea what can be done about that...
--- libgomp/testsuite/libgomp.c++/target-7.C.jj 2015-07-22 11:36:53.042867520
+0200
+++ libgomp/testsuite/libgomp.c++/target-7.C 2015-07-22 11:32:00.000000000
+0200
@@ -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);
+}
--- libgomp/testsuite/libgomp.c/target-15.c.jj 2015-07-22 11:37:11.655612690
+0200
+++ libgomp/testsuite/libgomp.c/target-15.c 2015-07-22 11:38:54.590203394
+0200
@@ -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;
+}
--- libgomp/target.c.jj 2015-07-21 09:07:23.690851224 +0200
+++ libgomp/target.c 2015-07-22 21:12:22.438213557 +0200
@@ -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
--- libgomp/libgomp.h.jj 2015-07-15 13:00:32.000000000 +0200
+++ libgomp/libgomp.h 2015-07-22 21:09:39.023307107 +0200
@@ -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;
};
--- include/gomp-constants.h.jj 2015-07-21 09:07:23.689851239 +0200
+++ include/gomp-constants.h 2015-07-21 15:01:05.384829637 +0200
@@ -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) \
--- gcc/cp/parser.c.jj 2015-07-21 09:06:42.000000000 +0200
+++ gcc/cp/parser.c 2015-07-21 12:55:03.966656803 +0200
@@ -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),
--- gcc/cp/semantics.c.jj 2015-07-17 13:59:27.000000000 +0200
+++ gcc/cp/semantics.c 2015-07-22 13:01:26.296499686 +0200
@@ -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)
--- gcc/tree.h.jj 2015-07-16 17:56:41.000000000 +0200
+++ gcc/tree.h 2015-07-21 16:40:35.759401307 +0200
@@ -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
--- gcc/gimplify.c.jj 2015-07-16 17:56:41.000000000 +0200
+++ gcc/gimplify.c 2015-07-22 18:27:58.545933111 +0200
@@ -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
--- gcc/c/c-tree.h.jj 2015-07-01 12:50:49.000000000 +0200
+++ gcc/c/c-tree.h 2015-07-22 12:47:49.185826677 +0200
@@ -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);
--- gcc/c/c-typeck.c.jj 2015-07-17 13:06:58.000000000 +0200
+++ gcc/c/c-typeck.c 2015-07-22 13:00:21.130399057 +0200
@@ -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)))
{
--- gcc/c/c-parser.c.jj 2015-07-21 09:06:42.000000000 +0200
+++ gcc/c/c-parser.c 2015-07-22 12:28:35.987814464 +0200
@@ -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);
}
--- gcc/tree-core.h.jj 2015-07-17 09:30:44.000000000 +0200
+++ gcc/tree-core.h 2015-07-21 16:28:48.524156167 +0200
@@ -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;
--- gcc/omp-low.c.jj 2015-07-21 09:07:23.000000000 +0200
+++ gcc/omp-low.c 2015-07-22 13:30:17.507589880 +0200
@@ -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);
}
--- gcc/tree-pretty-print.c.jj 2015-07-21 09:06:42.000000000 +0200
+++ gcc/tree-pretty-print.c 2015-07-22 13:53:51.406065024 +0200
@@ -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)
Jakub
