Patchwork [gomp4] Further accel fixes

login
register
mail settings
Submitter Jakub Jelinek
Date Sept. 6, 2013, 2:57 p.m.
Message ID <20130906145731.GS23437@tucnak.redhat.com>
Download mbox | patch
Permalink /patch/273215/
State New
Headers show

Comments

Jakub Jelinek - Sept. 6, 2013, 2:57 p.m.
Hi!

This fixes mainly VLA handling in target{, data, update} constructs,
but also deals with field alignments in the structure and field order.
Committed to gomp-4_0-branch.

2013-09-06  Jakub Jelinek  <jakub@redhat.com>

	* omp-low.c (scan_sharing_clauses): Handle VLAs in
	OMP_CLAUSE_{MAP,TO,FROM}.  Set DECL_ALIGN (field) before
	calling insert_field_into_struct.
	(scan_omp_target): Reverse TYPE_FIELDS, verify that
	all field alignments are the same.
	(lower_omp_target): Use maybe_lookup_field instead of
	lookup_sfield to check if field is present.  Handle VLAs.
	* tree-pretty-print.c (dump_omp_clause): Only check
	OMP_CLAUSE_MAP_KIND on OMP_CLAUSE_MAP clauses.
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_TO_ONLY.
	(omp_firstprivatize_variable, omp_add_variable,
	gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Handle
	VLAs in OMP_CLAUSE_{MAP,TO,FROM}.
libgomp/
	* testsuite/libgomp.c/target-2.c: New test.
	* testsuite/libgomp.c++/target-3.C: New test.


	Jakub

Patch

--- gcc/omp-low.c.jj	2013-09-05 17:11:14.000000000 +0200
+++ gcc/omp-low.c	2013-09-06 16:15:16.367638718 +0200
@@ -1574,10 +1574,24 @@  scan_sharing_clauses (tree clauses, omp_
 	    }
 	  if (DECL_P (decl))
 	    {
-	      install_var_field (decl, true, 3, ctx);
-	      if (gimple_omp_target_kind (ctx->stmt)
-		  == GF_OMP_TARGET_KIND_REGION)
-		install_var_local (decl, ctx);
+	      if (DECL_SIZE (decl)
+		  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+		{
+		  tree decl2 = DECL_VALUE_EXPR (decl);
+		  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+		  decl2 = TREE_OPERAND (decl2, 0);
+		  gcc_assert (DECL_P (decl2));
+		  install_var_field (decl2, true, 3, ctx);
+		  install_var_local (decl2, ctx);
+		  install_var_local (decl, ctx);
+		}
+	      else
+		{
+		  install_var_field (decl, true, 3, ctx);
+		  if (gimple_omp_target_kind (ctx->stmt)
+		      == GF_OMP_TARGET_KIND_REGION)
+		    install_var_local (decl, ctx);
+		}
 	    }
 	  else
 	    {
@@ -1600,6 +1614,7 @@  scan_sharing_clauses (tree clauses, omp_
 		  tree field
 		    = build_decl (OMP_CLAUSE_LOCATION (c),
 				  FIELD_DECL, NULL_TREE, ptr_type_node);
+		  DECL_ALIGN (field) = TYPE_ALIGN (ptr_type_node);
 		  insert_field_into_struct (ctx->record_type, field);
 		  splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
 				     (splay_tree_value) field);
@@ -1684,6 +1699,16 @@  scan_sharing_clauses (tree clauses, omp_
 		  TREE_TYPE (new_decl)
 		    = remap_type (TREE_TYPE (decl), &ctx->cb);
 		}
+	      else 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));
+		  fixup_remapped_decl (decl2, ctx, false);
+		  fixup_remapped_decl (decl, ctx, true);
+		}
 	      else
 		fixup_remapped_decl (decl, ctx, false);
 	    }
@@ -2126,6 +2151,16 @@  scan_omp_target (gimple stmt, omp_contex
     ctx->record_type = ctx->receiver_decl = NULL;
   else
     {
+      TYPE_FIELDS (ctx->record_type)
+	= nreverse (TYPE_FIELDS (ctx->record_type));
+#ifdef ENABLE_CHECKING
+      tree field;
+      unsigned int align = DECL_ALIGN (TYPE_FIELDS (ctx->record_type));
+      for (field = TYPE_FIELDS (ctx->record_type);
+	   field;
+	   field = DECL_CHAIN (field))
+	gcc_assert (DECL_ALIGN (field) == align);
+#endif
       layout_type (ctx->record_type);
       if (kind == GF_OMP_TARGET_KIND_REGION)
 	fixup_child_record_type (ctx);
@@ -9201,7 +9236,18 @@  lower_omp_target (gimple_stmt_iterator *
 	      map_cnt++;
 	    continue;
 	  }
-	if (!lookup_sfield (var, ctx))
+
+	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;
+	  }
+
+	if (!maybe_lookup_field (var, ctx))
 	  continue;
 
 	if (kind == GF_OMP_TARGET_KIND_REGION)
@@ -9293,8 +9339,20 @@  lower_omp_target (gimple_stmt_iterator *
 		    nc = NULL_TREE;
 		  }
 	      }
-	    else if (!lookup_sfield (ovar, ctx))
-	      continue;
+	    else
+	      {
+		if (DECL_SIZE (ovar)
+		    && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
+		  {
+		    tree ovar2 = DECL_VALUE_EXPR (ovar);
+		    gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
+		    ovar2 = TREE_OPERAND (ovar2, 0);
+		    gcc_assert (DECL_P (ovar2));
+		    ovar = ovar2;
+		  }
+		if (!maybe_lookup_field (ovar, ctx))
+		  continue;
+	      }
 
 	    if (nc)
 	      {
--- gcc/tree-pretty-print.c.jj	2013-08-27 22:18:05.000000000 +0200
+++ gcc/tree-pretty-print.c	2013-09-06 15:36:48.481578926 +0200
@@ -508,7 +508,8 @@  dump_omp_clause (pretty_printer *buffer,
      print_clause_size:
       if (OMP_CLAUSE_SIZE (clause))
 	{
-	  if (OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER)
+	  if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+	      && OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER)
 	    pp_string (buffer, " [pointer assign, bias: ");
 	  else
 	    pp_string (buffer, " [len: ");
--- gcc/gimplify.c.jj	2013-09-05 14:45:48.000000000 +0200
+++ gcc/gimplify.c	2013-09-06 14:47:13.153081283 +0200
@@ -61,6 +61,7 @@  enum gimplify_omp_var_data
   GOVD_PRIVATE_OUTER_REF = 1024,
   GOVD_LINEAR = 2048,
   GOVD_ALIGNED = 4096,
+  GOVD_MAP_TO_ONLY = 8192,
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -5740,11 +5741,16 @@  omp_firstprivatize_variable (struct gimp
 	{
 	  if (n->value & GOVD_SHARED)
 	    n->value = GOVD_FIRSTPRIVATE | (n->value & GOVD_SEEN);
+	  else if (n->value & GOVD_MAP)
+	    n->value |= GOVD_MAP_TO_ONLY;
 	  else
 	    return;
 	}
+      else if (ctx->region_type == ORT_TARGET)
+	omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
       else if (ctx->region_type != ORT_WORKSHARE
-	       && ctx->region_type != ORT_SIMD)
+	       && ctx->region_type != ORT_SIMD
+	       && ctx->region_type != ORT_TARGET_DATA)
 	omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
 
       ctx = ctx->outer_context;
@@ -5847,16 +5853,15 @@  omp_add_variable (struct gimplify_omp_ct
      the parameters of the type.  */
   if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
     {
-      /* To be handled later.  */
-      gcc_assert ((flags & GOVD_MAP) == 0);
-
       /* Add the pointer replacement variable as PRIVATE if the variable
 	 replacement is private, else FIRSTPRIVATE since we'll need the
 	 address of the original variable either for SHARED, or for the
 	 copy into or out of the context.  */
       if (!(flags & GOVD_LOCAL))
 	{
-	  nflags = flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
+	  nflags = flags & GOVD_MAP
+		   ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT
+		   : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
 	  nflags |= flags & GOVD_SEEN;
 	  t = DECL_VALUE_EXPR (decl);
 	  gcc_assert (TREE_CODE (t) == INDIRECT_REF);
@@ -5885,15 +5890,13 @@  omp_add_variable (struct gimplify_omp_ct
 	 For local variables TYPE_SIZE_UNIT might not be gimplified yet,
 	 in this case omp_notice_variable will be called later
 	 on when it is gimplified.  */
-      else if (! (flags & GOVD_LOCAL)
+      else if (! (flags & (GOVD_LOCAL | GOVD_MAP))
 	       && DECL_P (TYPE_SIZE_UNIT (TREE_TYPE (decl))))
 	omp_notice_variable (ctx, TYPE_SIZE_UNIT (TREE_TYPE (decl)), true);
     }
-  else if (lang_hooks.decls.omp_privatize_by_reference (decl))
+  else if ((flags & GOVD_MAP) == 0
+	   && lang_hooks.decls.omp_privatize_by_reference (decl))
     {
-      /* To be handled later.  */
-      gcc_assert ((flags & GOVD_MAP) == 0);
-
       gcc_assert ((flags & GOVD_LOCAL) == 0);
       omp_firstprivatize_type_sizes (ctx, TREE_TYPE (decl));
 
@@ -6562,7 +6565,35 @@  gimplify_adjust_omp_clauses_1 (splay_tre
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
     OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
   else if (code == OMP_CLAUSE_MAP)
-    OMP_CLAUSE_MAP_KIND (clause) = OMP_CLAUSE_MAP_TOFROM;
+    {
+      OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
+				     ? OMP_CLAUSE_MAP_TO
+				     : OMP_CLAUSE_MAP_TOFROM;
+      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));
+	  tree mem = build_simple_mem_ref (decl2);
+	  OMP_CLAUSE_DECL (clause) = mem;
+	  OMP_CLAUSE_SIZE (clause) = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+	  if (gimplify_omp_ctxp->outer_context)
+	    {
+	      struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp->outer_context;
+	      omp_notice_variable (ctx, decl2, true);
+	      omp_notice_variable (ctx, OMP_CLAUSE_SIZE (clause), true);
+	    }
+	  tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_DECL (nc) = decl;
+	  OMP_CLAUSE_SIZE (nc) = size_zero_node;
+	  OMP_CLAUSE_MAP_KIND (nc) = OMP_CLAUSE_MAP_POINTER;
+	  OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
+	  OMP_CLAUSE_CHAIN (clause) = nc;
+	}
+    }
   *list_p = clause;
   lang_hooks.decls.omp_finish_clause (clause);
 
@@ -6687,6 +6718,56 @@  gimplify_adjust_omp_clauses (tree *list_
 	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
 	  if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
 	    remove = true;
+	  else if (DECL_SIZE (decl)
+		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
+		   && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER)
+	    {
+	      tree decl2 = DECL_VALUE_EXPR (decl);
+	      gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
+	      decl2 = TREE_OPERAND (decl2, 0);
+	      gcc_assert (DECL_P (decl2));
+	      tree mem = build_simple_mem_ref (decl2);
+	      OMP_CLAUSE_DECL (c) = mem;
+	      OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+	      if (ctx->outer_context)
+		{
+		  omp_notice_variable (ctx->outer_context, decl2, true);
+		  omp_notice_variable (ctx->outer_context,
+				       OMP_CLAUSE_SIZE (c), true);
+		}
+	      tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					  OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_DECL (nc) = decl;
+	      OMP_CLAUSE_SIZE (nc) = size_zero_node;
+	      OMP_CLAUSE_MAP_KIND (nc) = OMP_CLAUSE_MAP_POINTER;
+	      OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
+	      OMP_CLAUSE_CHAIN (c) = nc;
+	      c = nc;
+	    }
+	  break;
+
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	  decl = OMP_CLAUSE_DECL (c);
+	  if (!DECL_P (decl))
+	    break;
+	  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));
+	      tree mem = build_simple_mem_ref (decl2);
+	      OMP_CLAUSE_DECL (c) = mem;
+	      OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+	      if (ctx->outer_context)
+		{
+		  omp_notice_variable (ctx->outer_context, decl2, true);
+		  omp_notice_variable (ctx->outer_context,
+				       OMP_CLAUSE_SIZE (c), true);
+		}
+	    }
 	  break;
 
 	case OMP_CLAUSE_REDUCTION:
@@ -6708,8 +6789,6 @@  gimplify_adjust_omp_clauses (tree *list_
 	case OMP_CLAUSE_MERGEABLE:
 	case OMP_CLAUSE_PROC_BIND:
 	case OMP_CLAUSE_SAFELEN:
-	case OMP_CLAUSE_TO:
-	case OMP_CLAUSE_FROM:
 	case OMP_CLAUSE_DEPEND:
 	  break;
 
--- libgomp/testsuite/libgomp.c/target-2.c.jj	2013-09-06 16:24:29.213769868 +0200
+++ libgomp/testsuite/libgomp.c/target-2.c	2013-09-06 16:24:05.000000000 +0200
@@ -0,0 +1,88 @@ 
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+
+void
+fn1 (double *x, double *y, int z)
+{
+  int i;
+  for (i = 0; i < z; i++)
+    {
+      x[i] = i & 31;
+      y[i] = (i & 63) - 30;
+    }
+}
+
+double
+fn2 (int x)
+{
+  double s = 0;
+  double b[3 * x], c[3 * x], d[3 * x], e[3 * x];
+  int i;
+  fn1 (b, c, x);
+  fn1 (e, d + x, x);
+  #pragma omp target map(to: b, c[:x], d[x:x], e)
+    #pragma omp parallel for reduction(+:s)
+      for (i = 0; i < x; i++)
+	s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
+  return s;
+}
+
+double
+fn3 (int x)
+{
+  double s = 0;
+  double b[3 * x], c[3 * x], d[3 * x], e[3 * x];
+  int i;
+  fn1 (b, c, x);
+  fn1 (e, d, x);
+  #pragma omp target
+    #pragma omp parallel for reduction(+:s)
+      for (i = 0; i < x; i++)
+	s += b[i] * c[i] + d[i];
+  return s;
+}
+
+double
+fn4 (int x)
+{
+  double s = 0;
+  double b[3 * x], c[3 * x], d[3 * x], e[3 * x];
+  int i;
+  fn1 (b, c, x);
+  fn1 (e, d + x, x);
+  #pragma omp target data map(from: b, c[:x], d[x:x], e)
+    {
+      #pragma omp target update to(b, c[:x], d[x:x], e)
+      #pragma omp target map(c[:x], d[x:x])
+	#pragma omp parallel for reduction(+:s)
+	  for (i = 0; i < x; i++)
+	    {
+	      s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c);
+	      b[i] = i + 0.5;
+	      c[i] = 0.5 - i;
+	      d[x + i] = 0.5 * i;
+	    }
+    }
+  for (i = 0; i < x; i++)
+    if (b[i] != i + 0.5 || c[i] != 0.5 - i || d[x + i] != 0.5 * i)
+      abort ();
+  return s;
+}
+
+int
+main ()
+{
+  double a = fn2 (128);
+  if (a != 14080.0)
+    abort ();
+  double b = fn3 (128);
+  if (a != b)
+    abort ();
+  double c = fn4 (256);
+  if (c != 28160.0)
+    abort ();
+  return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-3.C.jj	2013-09-06 16:24:40.674710426 +0200
+++ libgomp/testsuite/libgomp.c++/target-3.C	2013-09-06 16:24:44.594690089 +0200
@@ -0,0 +1 @@ 
+#include "../libgomp.c/target-2.c"