Re: [gomp4.1] Depend clause support for offloading

[Jakub Jelinek <jakub at redhat dot com>]

On Thu, Sep 03, 2015 at 04:16:50PM +0200, Jakub Jelinek wrote:
> On Wed, Sep 02, 2015 at 05:58:54PM +0200, Jakub Jelinek wrote:
> > Here is the start of the async offloading support I've talked about,
> > but nowait is not supported on the library side yet, only depend clause
> > (and for that I haven't added a testcase yet).
> 
> Added testcase revealed two (small) issues, here is a fix for that together
> with the testcase.

There has been a bug in the testcase (missing map(from:err) in 3 places),
which hid a problem that on target constructs with depend clause (what about
just nowait?) we have to avoid using GOMP_FIRSTPRIVATE_INT or copy value
into temporary and take temporary's address for GOMP_FIRSTPRIVATE unless
we can prove other tasks can't modify the value while waiting for
dependencies (if it is addressable or shared with other threads/tasks,
then we have to use GOMP_FIRSTPRIVATE with address of the real variable,
so that if other tasks change it, we pick up the right values).

2015-09-04  Jakub Jelinek  <jakub@redhat.com>

	* omp-low.c (lower_omp_target): If target has depend
	clauses, avoid using GOMP_MAP_FIRSTPRIVATE_INT unless the
	var is non-addressable and private in the current task.
	Even for GOMP_MAP_FIRSTPRIVATE, if the var is non-addressable,
	but shared or threadprivate, take address of the shared var
	rather than initializing a temporary with the current value.

	* testsuite/libgomp.c/target-25.c (main): Add missing
	map(from: err) clauses to target constructs.

--- gcc/omp-low.c.jj	2015-09-03 16:36:31.000000000 +0200
+++ gcc/omp-low.c	2015-09-04 11:34:45.512416693 +0200
@@ -13236,6 +13236,7 @@ lower_omp_target (gimple_stmt_iterator *
   location_t loc = gimple_location (stmt);
   bool offloaded, data_region;
   unsigned int map_cnt = 0;
+  bool has_depend = false;
 
   offloaded = is_gimple_omp_offloaded (stmt);
   switch (gimple_omp_target_kind (stmt))
@@ -13268,6 +13269,7 @@ lower_omp_target (gimple_stmt_iterator *
       dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
       lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt),
 			    &dep_ilist, &dep_olist);
+      has_depend = true;
     }
 
   tgt_bind = NULL;
@@ -13719,9 +13721,44 @@ lower_omp_target (gimple_stmt_iterator *
 	    type = TREE_TYPE (ovar);
 	    if (is_reference (ovar))
 	      type = TREE_TYPE (type);
+	    bool use_firstprivate_int, force_addr;
+	    use_firstprivate_int = false;
+	    force_addr = false;
 	    if ((INTEGRAL_TYPE_P (type)
-		  && TYPE_PRECISION (type) <= POINTER_SIZE)
+		 && TYPE_PRECISION (type) <= POINTER_SIZE)
 		|| TREE_CODE (type) == POINTER_TYPE)
+	      use_firstprivate_int = true;
+	    if (has_depend)
+	      {
+		if (is_reference (var))
+		  use_firstprivate_int = false;
+		else if (is_gimple_reg (var))
+		  {
+		    if (DECL_HAS_VALUE_EXPR_P (var))
+		      {
+			tree v = get_base_address (var);
+			if (DECL_P (v) && TREE_ADDRESSABLE (v))
+			  {
+			    use_firstprivate_int = false;
+			    force_addr = true;
+			  }
+			else
+			  switch (TREE_CODE (v))
+			    {
+			    case INDIRECT_REF:
+			    case MEM_REF:
+			      use_firstprivate_int = false;
+			      force_addr = true;
+			      break;
+			    default:
+			      break;
+			    }
+		      }
+		  }
+		else
+		  use_firstprivate_int = false;
+	      }
+	    if (use_firstprivate_int)
 	      {
 		tkind = GOMP_MAP_FIRSTPRIVATE_INT;
 		tree t = var;
@@ -13734,7 +13771,7 @@ lower_omp_target (gimple_stmt_iterator *
 	      }
 	    else if (is_reference (var))
 	      gimplify_assign (x, var, &ilist);
-	    else if (is_gimple_reg (var))
+	    else if (!force_addr && is_gimple_reg (var))
 	      {
 		tree avar = create_tmp_var (TREE_TYPE (var));
 		mark_addressable (avar);
@@ -13867,9 +13904,40 @@ lower_omp_target (gimple_stmt_iterator *
 		type = TREE_TYPE (var);
 		if (is_reference (var))
 		  type = TREE_TYPE (type);
+		bool use_firstprivate_int;
+		use_firstprivate_int = false;
 		if ((INTEGRAL_TYPE_P (type)
 		     && TYPE_PRECISION (type) <= POINTER_SIZE)
 		    || TREE_CODE (type) == POINTER_TYPE)
+		  use_firstprivate_int = true;
+		if (has_depend)
+		  {
+		    tree v = lookup_decl_in_outer_ctx (var, ctx);
+		    if (is_reference (v))
+		      use_firstprivate_int = false;
+		    else if (is_gimple_reg (v))
+		      {
+			if (DECL_HAS_VALUE_EXPR_P (v))
+			  {
+			    v = get_base_address (v);
+			    if (DECL_P (v) && TREE_ADDRESSABLE (v))
+			      use_firstprivate_int = false;
+			    else
+			      switch (TREE_CODE (v))
+				{
+				case INDIRECT_REF:
+				case MEM_REF:
+				  use_firstprivate_int = false;
+				  break;
+				default:
+				  break;
+				}
+			  }
+		      }
+		    else
+		      use_firstprivate_int = false;
+		  }
+		if (use_firstprivate_int)
 		  {
 		    x = build_receiver_ref (var, false, ctx);
 		    if (TREE_CODE (type) != POINTER_TYPE)
--- libgomp/testsuite/libgomp.c/target-25.c.jj	2015-09-04 10:41:52.371881507 +0200
+++ libgomp/testsuite/libgomp.c/target-25.c	2015-09-04 10:39:16.000000000 +0200
@@ -23,7 +23,7 @@ main ()
       usleep (7000);
       z = 3;
     }
-    #pragma omp target map(tofrom: x) firstprivate (y) depend(inout: x, z)
+    #pragma omp target map(tofrom: x) map(from: err) firstprivate (y) depend(inout: x, z)
     err = (x != 1 || y != 2 || z != 3);
     if (err)
       abort ();
@@ -44,7 +44,7 @@ main ()
     }
     #pragma omp target enter data nowait map (to: w)
     #pragma omp target enter data depend (inout: x, z) map (to: x, y, z)
-    #pragma omp target map (alloc: x, y, z)
+    #pragma omp target map (alloc: x, y, z) map(from: err)
     {
       err = (x != 4 || y != 5 || z != 6);
       x = 7;
@@ -54,7 +54,7 @@ main ()
     if (err)
       abort ();
     #pragma omp taskwait
-    #pragma omp target map (alloc: w)
+    #pragma omp target map (alloc: w) map(from: err)
     {
       err = w != 7;
       w = 17;


	Jakub