On Fri, Jul 17, 2015 at 06:43:06PM +0200, Jakub Jelinek wrote:
> > BTW, do you plan to remove GOMP_MAP_POINTER mappings from array sections?
> > The enter/exit patch for libgomp depends on this change.
>
> My current plan (for Monday and onwards) is to first implement firstprivate
> on target construct, once that works hack on the GOMP_MAP_POINTER
> replacement, and then rewrite the gimplification rules for target construct
> for the new 2.15.5 rules (so that this one does not really break all the
> target tests we need the first two working somehow).
Ok, so here is the first part of that, GOMP_MAP_FIRSTPRIVATE support as a
way to support firstprivate/is_device_ptr clauses on target construct (and
private
clause too, though that is compiler only change).
firstprivate VLAs aren't supported yet, but that will be a compiler only
change.
I'll commit this patch tomorrow.
2015-07-20 Jakub Jelinek <ja...@redhat.com>
gcc/
* omp-low.c (scan_sharing_clauses): Handle firstprivate
and is_device_ptr clauses on target region.
(lower_omp_target): Handle OMP_CLAUSE_FIRSTPRIVATE,
OMP_CLAUSE_IS_DEVICE_PTR and OMP_CLAUSE_PRIVATE.
include/
* gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_FIRSTPRIVATE.
libgomp/
* target.c (gomp_map_vars): Handle GOMP_MAP_FIRSTPRIVATE.
* testsuite/libgomp.c/target-13.c: New test.
* testsuite/libgomp.c/target-14.c: New test.
* testsuite/libgomp.c++/target-5.C: New test.
* testsuite/libgomp.c++/target-6.C: New test.
--- gcc/omp-low.c.jj 2015-07-16 18:09:25.000000000 +0200
+++ gcc/omp-low.c 2015-07-20 17:43:33.271401254 +0200
@@ -1930,6 +1930,10 @@ 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)
@@ -12929,6 +12933,21 @@ lower_omp_target (gimple_stmt_iterator *
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
map_cnt++;
+ 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);
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ break;
}
if (offloaded)
@@ -12994,7 +13013,8 @@ lower_omp_target (gimple_stmt_iterator *
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
- tree ovar, nc;
+ tree ovar, nc, s, purpose, var, x;
+ unsigned int talign;
default:
break;
@@ -13037,13 +13057,13 @@ lower_omp_target (gimple_stmt_iterator *
continue;
}
- unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+ talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
if (nc)
{
- tree var = lookup_decl_in_outer_ctx (ovar, ctx);
- tree x = build_sender_ref (ovar, ctx);
+ var = lookup_decl_in_outer_ctx (ovar, ctx);
+ x = build_sender_ref (ovar, ctx);
if (maybe_lookup_oacc_reduction (var, ctx))
{
gcc_checking_assert (offloaded
@@ -13092,11 +13112,11 @@ lower_omp_target (gimple_stmt_iterator *
gimplify_assign (x, var, &ilist);
}
}
- tree s = OMP_CLAUSE_SIZE (c);
+ s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
- tree purpose = size_int (map_idx++);
+ purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
@@ -13126,6 +13146,52 @@ lower_omp_target (gimple_stmt_iterator *
build_int_cstu (tkind_type, tkind));
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)));
+ else
+ talign = DECL_ALIGN_UNIT (ovar);
+ var = lookup_decl_in_outer_ctx (ovar, ctx);
+ x = build_sender_ref (ovar, ctx);
+ if (is_reference (var))
+ gimplify_assign (x, var, &ilist);
+ else if (is_gimple_reg (var))
+ {
+ tree avar = create_tmp_var (TREE_TYPE (var));
+ mark_addressable (avar);
+ gimplify_assign (avar, var, &ilist);
+ avar = build_fold_addr_expr (avar);
+ gimplify_assign (x, avar, &ilist);
+ }
+ else
+ {
+ var = build_fold_addr_expr (var);
+ gimplify_assign (x, var, &ilist);
+ }
+ if (is_reference (var))
+ s = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
+ else
+ s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
+ s = fold_convert (size_type_node, s);
+ purpose = size_int (map_idx++);
+ CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
+ 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);
+ tkind |= 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));
+ break;
}
gcc_assert (map_idx == map_cnt);
@@ -13173,6 +13239,57 @@ lower_omp_target (gimple_stmt_iterator *
if (offloaded)
{
+ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ tree var;
+ 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);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
+ break;
+ case OMP_CLAUSE_PRIVATE:
+ var = OMP_CLAUSE_DECL (c);
+ if (is_reference (var))
+ {
+ 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)));
+ 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);
+ 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);
+ }
+
+ x = fold_convert_loc (clause_loc, 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;
+ }
gimple_seq_add_seq (&new_body, tgt_body);
new_body = maybe_catch_exception (new_body);
}
--- include/gomp-constants.h.jj 2015-06-23 16:23:45.000000000 +0200
+++ include/gomp-constants.h 2015-07-20 12:27:58.103210763 +0200
@@ -72,6 +72,8 @@ enum gomp_map_kind
/* Is a device pointer. OMP_CLAUSE_SIZE for these is unused; is implicitly
POINTER_SIZE_UNITS. */
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),
/* Allocate. */
GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC),
/* ..., and copy to device. */
--- libgomp/target.c.jj 2015-07-15 13:00:32.000000000 +0200
+++ libgomp/target.c 2015-07-20 16:03:20.745931639 +0200
@@ -243,6 +243,7 @@ gomp_map_vars (struct gomp_device_descr
bool short_mapkind, bool is_target)
{
size_t i, tgt_align, tgt_size, not_found_cnt = 0;
+ bool has_firstprivate = false;
const int rshift = short_mapkind ? 8 : 3;
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
@@ -280,6 +281,18 @@ gomp_map_vars (struct gomp_device_descr
cur_node.host_end = cur_node.host_start + sizes[i];
else
cur_node.host_end = cur_node.host_start + sizeof (void *);
+ if ((kind & typemask) == GOMP_MAP_FIRSTPRIVATE)
+ {
+ tgt->list[i].key = NULL;
+
+ size_t align = (size_t) 1 << (kind >> rshift);
+ if (tgt_align < align)
+ tgt_align = align;
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += cur_node.host_end - cur_node.host_start;
+ has_firstprivate = true;
+ continue;
+ }
splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
if (n)
gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
@@ -348,9 +361,10 @@ gomp_map_vars (struct gomp_device_descr
tgt_size = mapnum * sizeof (void *);
tgt->array = NULL;
- if (not_found_cnt)
+ if (not_found_cnt || has_firstprivate)
{
- tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array));
+ if (not_found_cnt)
+ tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array));
splay_tree_node array = tgt->array;
size_t j;
@@ -360,6 +374,18 @@ gomp_map_vars (struct gomp_device_descr
int kind = get_kind (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL)
continue;
+ if ((kind & typemask) == GOMP_MAP_FIRSTPRIVATE)
+ {
+ size_t align = (size_t) 1 << (kind >> rshift);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt->list[i].offset = tgt_size;
+ size_t len = sizes[i];
+ devicep->host2dev_func (devicep->target_id,
+ (void *) (tgt->tgt_start + tgt_size),
+ (void *) hostaddrs[i], len);
+ tgt_size += len;
+ continue;
+ }
splay_tree_key k = &array->key;
k->host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
@@ -491,7 +517,13 @@ gomp_map_vars (struct gomp_device_descr
for (i = 0; i < mapnum; i++)
{
if (tgt->list[i].key == NULL)
- cur_node.tgt_offset = (uintptr_t) NULL;
+ {
+ if (hostaddrs[i] == NULL)
+ cur_node.tgt_offset = (uintptr_t) NULL;
+ 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;
--- libgomp/testsuite/libgomp.c/target-13.c.jj 2015-07-20 16:07:28.259375318
+0200
+++ libgomp/testsuite/libgomp.c/target-13.c 2015-07-20 16:26:05.828330031
+0200
@@ -0,0 +1,45 @@
+#ifdef __cplusplus
+extern "C"
+#else
+extern
+#endif
+void abort (void);
+struct S { int s, t; };
+
+void
+foo ()
+{
+ int x = 5, y = 6, err = 0;
+ struct S u = { 7, 8 }, v = { 9, 10 };
+ double s = 11.5, t = 12.5;
+ #pragma omp target private (x, u, s) firstprivate (y, v, t) map(from:err)
+ {
+ x = y;
+ u = v;
+ s = t;
+ err = (x != 6 || y != 6
+ || u.s != 9 || u.t != 10 || v.s != 9 || v.t != 10
+ || s != 12.5 || t != 12.5);
+ x += 1;
+ y += 2;
+ u.s += 3;
+ v.t += 4;
+ s += 2.5;
+ t += 3.0;
+ if (x != 7 || y != 8
+ || u.s != 12 || u.t != 10 || v.s != 9 || v.t != 14
+ || s != 15.0 || t != 15.5)
+ err = 1;
+ }
+ if (err || x != 5 || y != 6
+ || u.s != 7 || u.t != 8 || v.s != 9 || v.t != 10
+ || s != 11.5 || t != 12.5)
+ abort ();
+}
+
+int
+main ()
+{
+ foo ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/target-14.c.jj 2015-07-20 17:44:51.443299100
+0200
+++ libgomp/testsuite/libgomp.c/target-14.c 2015-07-20 17:49:20.745483458
+0200
@@ -0,0 +1,38 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int err;
+ void *p;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ p = omp_target_alloc (128 * sizeof (int), d);
+ if (p == NULL)
+ return 0;
+
+ #pragma omp target is_device_ptr (p) if (d >= 0) device (d >= 0 ? d : 0)
+ {
+ int i, *q = (int *) p;
+ for (i = 0; i < 128; i++)
+ q[i] = i + 7;
+ }
+ #pragma omp target is_device_ptr (p) if (d >= 0) device (d >= 0 ? d : 0)
map(from:err)
+ {
+ int i;
+ err = 0;
+ for (i = 0; i < 128; i++)
+ if (((int *) p)[i] != i + 7)
+ err = 1;
+ }
+ if (err)
+ abort ();
+
+ omp_target_free (p, d);
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-5.C.jj 2015-07-20 16:23:39.592423836
+0200
+++ libgomp/testsuite/libgomp.c++/target-5.C 2015-07-20 16:26:31.049968908
+0200
@@ -0,0 +1 @@
+#include "../libgomp.c/target-13.c"
--- libgomp/testsuite/libgomp.c++/target-6.C.jj 2015-07-20 16:26:44.196780672
+0200
+++ libgomp/testsuite/libgomp.c++/target-6.C 2015-07-20 17:36:18.357533147
+0200
@@ -0,0 +1,64 @@
+extern "C" void abort (void);
+struct S { int s, t; };
+
+void
+foo (int &x, int &y, S &u, S &v, double &s, double &t)
+{
+ int err = 0, i;
+ int a[y - 2], b[y - 2];
+ int (&c)[y - 2] = a, (&d)[y - 2] = b;
+ for (i = 0; i < y - 2; i++)
+ {
+ c[i] = i;
+ d[i] = 3 + i;
+ }
+ #pragma omp target private (x, u, s, c, i) firstprivate (y, v, t, d)
map(from:err)
+ {
+ x = y;
+ u = v;
+ s = t;
+ for (i = 0; i < y - 2; i++)
+ c[i] = d[i];
+ err = (x != 6 || y != 6
+ || u.s != 9 || u.t != 10 || v.s != 9 || v.t != 10
+ || s != 12.5 || t != 12.5);
+ for (i = 0; i < y - 2; i++)
+ if (d[i] != 3 + i || c[i] != 3 + i)
+ err = 1;
+ else
+ {
+ c[i] += 2 * i;
+ d[i] += i;
+ }
+ x += 1;
+ y += 2;
+ u.s += 3;
+ v.t += 4;
+ s += 2.5;
+ t += 3.0;
+ if (x != 7 || y != 8
+ || u.s != 12 || u.t != 10 || v.s != 9 || v.t != 14
+ || s != 15.0 || t != 15.5)
+ err = 1;
+ for (i = 0; i < y - 4; i++)
+ if (d[i] != 3 + 2 * i || c[i] != 3 + 3 * i)
+ err = 1;
+ }
+ if (err || x != 5 || y != 6
+ || u.s != 7 || u.t != 8 || v.s != 9 || v.t != 10
+ || s != 11.5 || t != 12.5)
+ abort ();
+ for (i = 0; i < y - 2; i++)
+ if (d[i] != 3 + i || c[i] != i)
+ abort ();
+}
+
+int
+main ()
+{
+ int x = 5, y = 6;
+ S u = { 7, 8 }, v = { 9, 10 };
+ double s = 11.5, t = 12.5;
+ foo (x, y, u, v, s, t);
+ return 0;
+}
Jakub
