[gomp4.1] Initial support for some OpenMP 4.1 construct parsing

Jakub Jelinek jakub@redhat.com
Mon Jul 20 16:18:00 GMT 2015


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  <jakub@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



More information about the Gcc-patches mailing list