* omp-low.c (lower_omp_target): Use GOMP_MAP_FIRSTPRIVATE_INT

	regardless whether there are depend clauses or not.

	* libgomp.h (struct gomp_target_task): Remove firstprivate_copies
	field.
	* target.c (gomp_target_fallback_firstprivate,
	gomp_target_unshare_firstprivate): Removed.
	(GOMP_target_ext): Copy firstprivate vars into gomp_allocaed memory
	before waiting for dependencies.
	(gomp_target_task_fn): Don't copy firstprivate vars here.
	* task.c (GOMP_PLUGIN_target_task_completion): Don't free
	firstprivate_copies here.
	(gomp_create_target_task): Don't initialize firstprivate_copies field.
	* testsuite/libgomp.c/target-25.c (main): Use map (to:) instead of
	explicit/implicit firstprivate.


git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@234894 138bc75d-0d04-0410-961f-82ee72b054a4

5 files changed, 55 insertions, 58 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 742f19052a4..b53dc6b7056 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,18 @@
+2016-04-12  Jakub Jelinek  <jakub@redhat.com>
+
+	* libgomp.h (struct gomp_target_task): Remove firstprivate_copies
+	field.
+	* target.c (gomp_target_fallback_firstprivate,
+	gomp_target_unshare_firstprivate): Removed.
+	(GOMP_target_ext): Copy firstprivate vars into gomp_allocaed memory
+	before waiting for dependencies.
+	(gomp_target_task_fn): Don't copy firstprivate vars here.
+	* task.c (GOMP_PLUGIN_target_task_completion): Don't free
+	firstprivate_copies here.
+	(gomp_create_target_task): Don't initialize firstprivate_copies field.
+	* testsuite/libgomp.c/target-25.c (main): Use map (to:) instead of
+	explicit/implicit firstprivate.
+
 2016-04-08  Cesar Philippidis  <cesar@codesourcery.com>
 
 	PR lto/70289
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 7108a6d0118..664e76b52d1 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -496,8 +496,6 @@ struct gomp_target_task
   struct target_mem_desc *tgt;
   struct gomp_task *task;
   struct gomp_team *team;
-  /* Copies of firstprivate mapped data for shared memory accelerators.  */
-  void *firstprivate_copies;
   /* Device-specific target arguments.  */
   void **args;
   void *hostaddrs[];
diff --git a/libgomp/target.c b/libgomp/target.c
index 96fe3d5eb0d..e2dd0e08997 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1372,47 +1372,6 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs,
       }
 }
 
-/* Host fallback with firstprivate map-type handling.  */
-
-static void
-gomp_target_fallback_firstprivate (void (*fn) (void *), size_t mapnum,
-				   void **hostaddrs, size_t *sizes,
-				   unsigned short *kinds)
-{
-  size_t tgt_align = 0, tgt_size = 0;
-  calculate_firstprivate_requirements (mapnum, sizes, kinds, &tgt_align,
-				       &tgt_size);
-  if (tgt_align)
-    {
-      char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
-      copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, tgt_align,
-			      tgt_size);
-    }
-  gomp_target_fallback (fn, hostaddrs);
-}
-
-/* Handle firstprivate map-type for shared memory devices and the host
-   fallback.  Return the pointer of firstprivate copies which has to be freed
-   after use.  */
-
-static void *
-gomp_target_unshare_firstprivate (size_t mapnum, void **hostaddrs,
-				  size_t *sizes, unsigned short *kinds)
-{
-  size_t tgt_align = 0, tgt_size = 0;
-  char *tgt = NULL;
-
-  calculate_firstprivate_requirements (mapnum, sizes, kinds, &tgt_align,
-				       &tgt_size);
-  if (tgt_align)
-    {
-      tgt = gomp_malloc (tgt_size + tgt_align - 1);
-      copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, tgt_align,
-			      tgt_size);
-    }
-  return tgt;
-}
-
 /* Helper function of GOMP_target{,_ext} routines.  */
 
 static void *
@@ -1504,6 +1463,8 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
 		 unsigned int flags, void **depend, void **args)
 {
   struct gomp_device_descr *devicep = resolve_device (device);
+  size_t tgt_align = 0, tgt_size = 0;
+  bool fpc_done = false;
 
   if (flags & GOMP_TARGET_FLAG_NOWAIT)
     {
@@ -1555,7 +1516,19 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
     {
       struct gomp_thread *thr = gomp_thread ();
       if (thr->task && thr->task->depend_hash)
-	gomp_task_maybe_wait_for_dependencies (depend);
+	{
+	  /* If we might need to wait, copy firstprivate now.  */
+	  calculate_firstprivate_requirements (mapnum, sizes, kinds,
+					       &tgt_align, &tgt_size);
+	  if (tgt_align)
+	    {
+	      char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
+	      copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds,
+				      tgt_align, tgt_size);
+	    }
+	  fpc_done = true;
+	  gomp_task_maybe_wait_for_dependencies (depend);
+	}
     }
 
   void *fn_addr;
@@ -1564,15 +1537,35 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
       || !(fn_addr = gomp_get_target_fn_addr (devicep, fn))
       || (devicep->can_run_func && !devicep->can_run_func (fn_addr)))
     {
-      gomp_target_fallback_firstprivate (fn, mapnum, hostaddrs, sizes, kinds);
+      if (!fpc_done)
+	{
+	  calculate_firstprivate_requirements (mapnum, sizes, kinds,
+					       &tgt_align, &tgt_size);
+	  if (tgt_align)
+	    {
+	      char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
+	      copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds,
+				      tgt_align, tgt_size);
+	    }
+	}
+      gomp_target_fallback (fn, hostaddrs);
       return;
     }
 
   struct target_mem_desc *tgt_vars;
-  void *fpc = NULL;
   if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
     {
-      fpc = gomp_target_unshare_firstprivate (mapnum, hostaddrs, sizes, kinds);
+      if (!fpc_done)
+	{
+	  calculate_firstprivate_requirements (mapnum, sizes, kinds,
+					       &tgt_align, &tgt_size);
+	  if (tgt_align)
+	    {
+	      char *tgt = gomp_alloca (tgt_size + tgt_align - 1);
+	      copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds,
+				      tgt_align, tgt_size);
+	    }
+	}
       tgt_vars = NULL;
     }
   else
@@ -1583,8 +1576,6 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
 		     args);
   if (tgt_vars)
     gomp_unmap_vars (tgt_vars, true);
-  else
-    free (fpc);
 }
 
 /* Host fallback for GOMP_target_data{,_ext} routines.  */
@@ -1891,9 +1882,7 @@ gomp_target_task_fn (void *data)
 	  || (devicep->can_run_func && !devicep->can_run_func (fn_addr)))
 	{
 	  ttask->state = GOMP_TARGET_TASK_FALLBACK;
-	  gomp_target_fallback_firstprivate (ttask->fn, ttask->mapnum,
-					     ttask->hostaddrs, ttask->sizes,
-					     ttask->kinds);
+	  gomp_target_fallback (ttask->fn, ttask->hostaddrs);
 	  return false;
 	}
 
@@ -1908,9 +1897,6 @@ gomp_target_task_fn (void *data)
       if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
 	{
 	  ttask->tgt = NULL;
-	  ttask->firstprivate_copies
-	    = gomp_target_unshare_firstprivate (ttask->mapnum, ttask->hostaddrs,
-						ttask->sizes, ttask->kinds);
 	  actual_arguments = ttask->hostaddrs;
 	}
       else
diff --git a/libgomp/task.c b/libgomp/task.c
index 38d4e9b413b..023663f43d5 100644
--- a/libgomp/task.c
+++ b/libgomp/task.c
@@ -582,7 +582,6 @@ GOMP_PLUGIN_target_task_completion (void *data)
       return;
     }
   ttask->state = GOMP_TARGET_TASK_FINISHED;
-  free (ttask->firstprivate_copies);
   gomp_target_task_completion (team, task);
   gomp_mutex_unlock (&team->task_lock);
 }
@@ -683,7 +682,6 @@ gomp_create_target_task (struct gomp_device_descr *devicep,
   ttask->state = state;
   ttask->task = task;
   ttask->team = team;
-  ttask->firstprivate_copies = NULL;
   task->fn = NULL;
   task->fn_data = ttask;
   task->final_task = 0;
diff --git a/libgomp/testsuite/libgomp.c/target-25.c b/libgomp/testsuite/libgomp.c/target-25.c
index aeb19aee510..09b8d52184a 100644
--- a/libgomp/testsuite/libgomp.c/target-25.c
+++ b/libgomp/testsuite/libgomp.c/target-25.c
@@ -23,7 +23,7 @@ main ()
       usleep (7000);
       z = 3;
     }
-    #pragma omp target map(tofrom: x) map(from: err) firstprivate (y) depend(inout: x, z)
+    #pragma omp target map(tofrom: x) map(from: err) map (to: y, z) depend(inout: x, z)
     err = (x != 1 || y != 2 || z != 3);
     if (err)
       abort ();