[gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data

Ilya Verbin iverbin@gmail.com
Mon Jun 15 19:54:00 GMT 2015


On Mon, Jun 15, 2015 at 18:25:28 +0200, Jakub Jelinek wrote:
> On Mon, Jun 15, 2015 at 07:18:27PM +0300, Ilya Verbin wrote:
> > On Mon, Jun 15, 2015 at 15:06:09 +0200, Jakub Jelinek wrote:
> > > On Mon, Jun 15, 2015 at 03:20:37PM +0300, Ilya Verbin wrote:
> > > > This patch introduces new versions of GOMP_target{,_data,_update} for OpenMP 4.1
> > > > with unsigned short for map kinds, but without new async arguments yet.
> > > 
> > > I think I'd prefer (for now) to suffix the functions with _41 instead of 1
> > > (and we'll see if we can come up with better names when async support is
> > > added).
> > 
> > OK.
> 
> Thanks.
> 
> > > Do we need to change GOMP_target_update though (at least right
> > > now)?  I mean, the construct only allows to and from clauses, not the map
> > > clause, and those don't really have an always modifier, nor release/delete
> > > semantics etc., so at least for now I think using the current
> > > GOMP_target_update should be ok.
> > 
> > I thought that it wouldn't look good, since without GOMP_target_update_41 we
> > will need to keep this obsolete parts:
> 
> I'd prefer to keep it for now, perhaps later on we'll switch to 16-bit kinds
> even for that, but better figure out first what to do with the async stuff,
> handle the enter/exit data correctly, change the library for OpenMP 4.1 to
> do the fully refcounted model.

Here is the new patch.  OK to commit?


gcc/
	* builtin-types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR): New.
	(BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): Remove.
	* omp-builtins.def (BUILT_IN_GOMP_TARGET): Replace GOMP_target with
	GOMP_target_41.
	(BUILT_IN_GOMP_TARGET_DATA): Replace GOMP_target_data with
	GOMP_target_data_41.
	(BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA): New.
	* omp-low.c (expand_omp_target): Use
	BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA for GF_OMP_TARGET_KIND_ENTER_DATA
	and GF_OMP_TARGET_KIND_EXIT_DATA.
	Do not pass obsolete pointer to new builtins.
	(lower_omp_target): Use unsigned short for map kinds, except
	BUILT_IN_GOMP_TARGET_UPDATE.
gcc/fortran/
	* types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR): New.
	(BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): Remove.
libgomp/
	* libgomp.map (GOMP_4.1): Add GOMP_target_41, GOMP_target_data_41,
	GOMP_target_enter_exit_data.
	* libgomp_g.h: Declare GOMP_target_41, GOMP_target_data_41,
	GOMP_target_enter_exit_data.
	* target.c (resolve_device): Call gomp_init_device here instead of
	GOMP_target*.
	(get_kind): Rename is_openacc to short_mapkind.
	(gomp_map_vars): Likewise.
	(gomp_unmap_vars): Likewise.
	(gomp_update): Likewise.
	(gomp_target_fallback): New static function.
	(gomp_get_target_fn_addr): New static function.
	(GOMP_target): Move host fallback and fn lookup to the new functions.
	(GOMP_target_41): New function.
	(gomp_target_data_fallback): New static function.
	(GOMP_target_data): Move host fallback to the new function.
	(GOMP_target_data_41): New function.
	(GOMP_target_update): Do not call gomp_init_device.
	(GOMP_target_enter_exit_data): New function.


diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 492ca63..870c957 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -526,6 +526,9 @@ DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_SIZE_VPTR_PTR_PTR_INT_INT, BT_BOOL, BT_SIZE,
 		     BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT)
 DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
 		     BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR,
+		     BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
+		     BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_7 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT,
 		     BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
@@ -534,9 +537,6 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL_ULL_ULL_ULL_ULLPTR_ULLPTR,
 		     BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
 		     BT_ULONGLONG, BT_ULONGLONG,
 		     BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
-DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
-		     BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
-		     BT_PTR, BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
 		     BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index c0d3989..a830235 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -189,6 +189,9 @@ DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_VPTR_PTR_I16_BOOL_INT_INT,
 		     BT_INT)
 DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_SIZE_VPTR_PTR_PTR_INT_INT, BT_BOOL, BT_SIZE,
 		     BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR,
+		     BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
+		     BT_PTR, BT_PTR)
 DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
 		     BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
 
@@ -199,9 +202,6 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL_ULL_ULL_ULL_ULLPTR_ULLPTR,
 		     BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
 		     BT_ULONGLONG, BT_ULONGLONG,
 		     BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
-DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
-		     BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
-		     BT_PTR, BT_PTR, BT_PTR)
 
 DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
 		     BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 749def4..470f038 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -262,14 +262,16 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_START, "GOMP_single_copy_start",
 		  BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end",
 		  BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target",
-		  BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR,
-		  ATTR_NOTHROW_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data",
-		  BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target_41",
+		  BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data_41",
+		  BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data",
 		  BT_FN_VOID, ATTR_NOTHROW_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
 		  BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
+		  "GOMP_target_enter_exit_data",
+		  BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
 		  BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e4f5566..3e27f8a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -10226,8 +10226,7 @@ expand_omp_target (struct omp_region *region)
       break;
     case GF_OMP_TARGET_KIND_ENTER_DATA:
     case GF_OMP_TARGET_KIND_EXIT_DATA:
-      /* FIXME */
-      start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+      start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA;
       break;
     case GF_OMP_TARGET_KIND_OACC_PARALLEL:
     case GF_OMP_TARGET_KIND_OACC_KERNELS:
@@ -10264,7 +10263,8 @@ expand_omp_target (struct omp_region *region)
 	 defined/used for the OpenMP target ones.  */
       gcc_checking_assert (start_ix == BUILT_IN_GOMP_TARGET
 			   || start_ix == BUILT_IN_GOMP_TARGET_DATA
-			   || start_ix == BUILT_IN_GOMP_TARGET_UPDATE);
+			   || start_ix == BUILT_IN_GOMP_TARGET_UPDATE
+			   || start_ix == BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA);
 
       device = OMP_CLAUSE_DEVICE_ID (c);
       clause_loc = OMP_CLAUSE_LOCATION (c);
@@ -10351,23 +10351,10 @@ expand_omp_target (struct omp_region *region)
   args.quick_push (device);
   if (offloaded)
     args.quick_push (build_fold_addr_expr (child_fn));
-  switch (start_ix)
-    {
-    case BUILT_IN_GOMP_TARGET:
-    case BUILT_IN_GOMP_TARGET_DATA:
-    case BUILT_IN_GOMP_TARGET_UPDATE:
-      /* This const void * is part of the current ABI, but we're not actually
-	 using it.  */
-      args.quick_push (build_zero_cst (ptr_type_node));
-      break;
-    case BUILT_IN_GOACC_DATA_START:
-    case BUILT_IN_GOACC_ENTER_EXIT_DATA:
-    case BUILT_IN_GOACC_PARALLEL:
-    case BUILT_IN_GOACC_UPDATE:
-      break;
-    default:
-      gcc_unreachable ();
-    }
+  /* This const void * is part of the current ABI, but we're not actually using
+     it.  */
+  if (start_ix == BUILT_IN_GOMP_TARGET_UPDATE)
+    args.quick_push (build_zero_cst (ptr_type_node));
   args.quick_push (t1);
   args.quick_push (t2);
   args.quick_push (t3);
@@ -10378,6 +10365,7 @@ expand_omp_target (struct omp_region *region)
     case BUILT_IN_GOMP_TARGET:
     case BUILT_IN_GOMP_TARGET_DATA:
     case BUILT_IN_GOMP_TARGET_UPDATE:
+    case BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA:
       break;
     case BUILT_IN_GOACC_PARALLEL:
       {
@@ -12633,7 +12621,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
       tree tkind_type;
       int talign_shift;
-      if (is_gimple_omp_oacc (stmt))
+      if (is_gimple_omp_oacc (stmt)
+	  || gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_UPDATE)
 	{
 	  tkind_type = short_unsigned_type_node;
 	  talign_shift = 8;
@@ -12782,9 +12771,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	      default:
 		gcc_unreachable ();
 	      }
-	    /* FIXME: Temporary hack.  */
-	    if (talign_shift == 3)
-	      tkind &= ~GOMP_MAP_FLAG_FORCE;
 	    gcc_checking_assert (tkind
 				 < (HOST_WIDE_INT_C (1U) << talign_shift));
 	    talign = ceil_log2 (talign);
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 36c0bb5..a77f1e3 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -242,6 +242,9 @@ GOMP_4.0.1 {
 
 GOMP_4.1 {
   global:
+	GOMP_target_41;
+	GOMP_target_data_41;
+	GOMP_target_enter_exit_data;
 	GOMP_taskloop;
 	GOMP_taskloop_ull;
 } GOMP_4.0.1;
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 5e88d45..ef7dc0d 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -208,11 +208,17 @@ extern void GOMP_single_copy_end (void *);
 
 extern void GOMP_target (int, void (*) (void *), const void *,
 			 size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_41 (int, void (*) (void *), size_t, void **, size_t *,
+			  unsigned short *);
 extern void GOMP_target_data (int, const void *,
 			      size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_data_41 (int, size_t, void **, size_t *,
+			       unsigned short *);
 extern void GOMP_target_end_data (void);
 extern void GOMP_target_update (int, const void *,
 				size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_enter_exit_data (int, size_t, void **, size_t *,
+					 unsigned short *);
 extern void GOMP_teams (unsigned int, unsigned int);
 
 /* oacc-parallel.c */
diff --git a/libgomp/target.c b/libgomp/target.c
index d8da783..218b1a4 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -132,6 +132,11 @@ resolve_device (int device_id)
   if (device_id < 0 || device_id >= gomp_get_num_devices ())
     return NULL;
 
+  gomp_mutex_lock (&devices[device_id].lock);
+  if (!devices[device_id].is_initialized)
+    gomp_init_device (&devices[device_id]);
+  gomp_mutex_unlock (&devices[device_id].lock);
+
   return &devices[device_id];
 }
 
@@ -157,20 +162,20 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn,
 }
 
 static int
-get_kind (bool is_openacc, void *kinds, int idx)
+get_kind (bool short_mapkind, void *kinds, int idx)
 {
-  return is_openacc ? ((unsigned short *) kinds)[idx]
-		    : ((unsigned char *) kinds)[idx];
+  return short_mapkind ? ((unsigned short *) kinds)[idx]
+		       : ((unsigned char *) kinds)[idx];
 }
 
 attribute_hidden struct target_mem_desc *
 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
-	       bool is_openacc, bool is_target)
+	       bool short_mapkind, bool is_target)
 {
   size_t i, tgt_align, tgt_size, not_found_cnt = 0;
-  const int rshift = is_openacc ? 8 : 3;
-  const int typemask = is_openacc ? 0xff : 0x7;
+  const int rshift = short_mapkind ? 8 : 3;
+  const int typemask = short_mapkind ? 0xff : 0x7;
   struct splay_tree_s *mem_map = &devicep->mem_map;
   struct splay_tree_key_s cur_node;
   struct target_mem_desc *tgt
@@ -195,7 +200,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 
   for (i = 0; i < mapnum; i++)
     {
-      int kind = get_kind (is_openacc, kinds, i);
+      int kind = get_kind (short_mapkind, kinds, i);
       if (hostaddrs[i] == NULL)
 	{
 	  tgt->list[i] = NULL;
@@ -226,7 +231,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	    {
 	      size_t j;
 	      for (j = i + 1; j < mapnum; j++)
-		if (!GOMP_MAP_POINTER_P (get_kind (is_openacc, kinds, j)
+		if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind, kinds, j)
 					 & typemask))
 		  break;
 		else if ((uintptr_t) hostaddrs[j] < cur_node.host_start
@@ -285,7 +290,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
       for (i = 0; i < mapnum; i++)
 	if (tgt->list[i] == NULL)
 	  {
-	    int kind = get_kind (is_openacc, kinds, i);
+	    int kind = get_kind (short_mapkind, kinds, i);
 	    if (hostaddrs[i] == NULL)
 	      continue;
 	    splay_tree_key k = &array->key;
@@ -394,7 +399,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 					    k->host_end - k->host_start);
 
 		    for (j = i + 1; j < mapnum; j++)
-		      if (!GOMP_MAP_POINTER_P (get_kind (is_openacc, kinds, j)
+		      if (!GOMP_MAP_POINTER_P (get_kind (short_mapkind, kinds,
+							 j)
 					       & typemask))
 			break;
 		      else if ((uintptr_t) hostaddrs[j] < k->host_start
@@ -613,11 +619,11 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
 
 static void
 gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
-	     size_t *sizes, void *kinds, bool is_openacc)
+	     size_t *sizes, void *kinds, bool short_mapkind)
 {
   size_t i;
   struct splay_tree_key_s cur_node;
-  const int typemask = is_openacc ? 0xff : 0x7;
+  const int typemask = short_mapkind ? 0xff : 0x7;
 
   if (!devicep)
     return;
@@ -634,7 +640,7 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 	splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
 	if (n)
 	  {
-	    int kind = get_kind (is_openacc, kinds, i);
+	    int kind = get_kind (short_mapkind, kinds, i);
 	    if (n->host_start > cur_node.host_start
 		|| n->host_end < cur_node.host_end)
 	      {
@@ -931,6 +937,47 @@ gomp_fini_device (struct gomp_device_descr *devicep)
   devicep->is_initialized = false;
 }
 
+/* Host fallback for GOMP_target{,_41} routines.  */
+
+static void
+gomp_target_fallback (void (*fn) (void *), void **hostaddrs)
+{
+  struct gomp_thread old_thr, *thr = gomp_thread ();
+  old_thr = *thr;
+  memset (thr, '\0', sizeof (*thr));
+  if (gomp_places_list)
+    {
+      thr->place = old_thr.place;
+      thr->ts.place_partition_len = gomp_places_list_len;
+    }
+  fn (hostaddrs);
+  gomp_free_thread (thr);
+  *thr = old_thr;
+}
+
+/* Helper function of GOMP_target{,_41} routines.  */
+
+static void *
+gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
+			 void (*host_fn) (void *))
+{
+  if (devicep->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC)
+    return (void *) host_fn;
+  else
+    {
+      gomp_mutex_lock (&devicep->lock);
+      struct splay_tree_key_s k;
+      k.host_start = (uintptr_t) host_fn;
+      k.host_end = k.host_start + 1;
+      splay_tree_key tgt_fn = splay_tree_lookup (&devicep->mem_map, &k);
+      gomp_mutex_unlock (&devicep->lock);
+      if (tgt_fn == NULL)
+	gomp_fatal ("Target function wasn't mapped");
+
+      return (void *) tgt_fn->tgt_offset;
+    }
+}
+
 /* Called when encountering a target directive.  If DEVICE
    is GOMP_DEVICE_ICV, it means use device-var ICV.  If it is
    GOMP_DEVICE_HOST_FALLBACK (or any value
@@ -950,50 +997,41 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
 
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
+    return gomp_target_fallback (fn, hostaddrs);
+
+  void *fn_addr = gomp_get_target_fn_addr (devicep, fn);
+
+  struct target_mem_desc *tgt_vars
+    = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
+		     true);
+  struct gomp_thread old_thr, *thr = gomp_thread ();
+  old_thr = *thr;
+  memset (thr, '\0', sizeof (*thr));
+  if (gomp_places_list)
     {
-      /* Host fallback.  */
-      struct gomp_thread old_thr, *thr = gomp_thread ();
-      old_thr = *thr;
-      memset (thr, '\0', sizeof (*thr));
-      if (gomp_places_list)
-	{
-	  thr->place = old_thr.place;
-	  thr->ts.place_partition_len = gomp_places_list_len;
-	}
-      fn (hostaddrs);
-      gomp_free_thread (thr);
-      *thr = old_thr;
-      return;
+      thr->place = old_thr.place;
+      thr->ts.place_partition_len = gomp_places_list_len;
     }
+  devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start);
+  gomp_free_thread (thr);
+  *thr = old_thr;
+  gomp_unmap_vars (tgt_vars, true);
+}
 
-  gomp_mutex_lock (&devicep->lock);
-  if (!devicep->is_initialized)
-    gomp_init_device (devicep);
-  gomp_mutex_unlock (&devicep->lock);
-
-  void *fn_addr;
+void
+GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
+		void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+  struct gomp_device_descr *devicep = resolve_device (device);
 
-  if (devicep->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC)
-    fn_addr = (void *) fn;
-  else
-    {
-      gomp_mutex_lock (&devicep->lock);
-      struct splay_tree_key_s k;
-      k.host_start = (uintptr_t) fn;
-      k.host_end = k.host_start + 1;
-      splay_tree_key tgt_fn = splay_tree_lookup (&devicep->mem_map, &k);
-      if (tgt_fn == NULL)
-	{
-	  gomp_mutex_unlock (&devicep->lock);
-	  gomp_fatal ("Target function wasn't mapped");
-	}
-      gomp_mutex_unlock (&devicep->lock);
+  if (devicep == NULL
+      || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
+    return gomp_target_fallback (fn, hostaddrs);
 
-      fn_addr = (void *) tgt_fn->tgt_offset;
-    }
+  void *fn_addr = gomp_get_target_fn_addr (devicep, fn);
 
   struct target_mem_desc *tgt_vars
-    = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
+    = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
 		     true);
   struct gomp_thread old_thr, *thr = gomp_thread ();
   old_thr = *thr;
@@ -1009,6 +1047,25 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
   gomp_unmap_vars (tgt_vars, true);
 }
 
+/* Host fallback for GOMP_target_data{,_41} routines.  */
+
+static void
+gomp_target_data_fallback (void)
+{
+  struct gomp_task_icv *icv = gomp_icv (false);
+  if (icv->target_data)
+    {
+      /* Even when doing a host fallback, if there are any active
+         #pragma omp target data constructs, need to remember the
+         new #pragma omp target data, otherwise GOMP_target_end_data
+         would get out of sync.  */
+      struct target_mem_desc *tgt
+	= gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, false);
+      tgt->prev = icv->target_data;
+      icv->target_data = tgt;
+    }
+}
+
 void
 GOMP_target_data (int device, const void *unused, size_t mapnum,
 		  void **hostaddrs, size_t *sizes, unsigned char *kinds)
@@ -1017,27 +1074,7 @@ GOMP_target_data (int device, const void *unused, size_t mapnum,
 
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
-    {
-      /* Host fallback.  */
-      struct gomp_task_icv *icv = gomp_icv (false);
-      if (icv->target_data)
-	{
-	  /* Even when doing a host fallback, if there are any active
-	     #pragma omp target data constructs, need to remember the
-	     new #pragma omp target data, otherwise GOMP_target_end_data
-	     would get out of sync.  */
-	  struct target_mem_desc *tgt
-	    = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, false);
-	  tgt->prev = icv->target_data;
-	  icv->target_data = tgt;
-	}
-      return;
-    }
-
-  gomp_mutex_lock (&devicep->lock);
-  if (!devicep->is_initialized)
-    gomp_init_device (devicep);
-  gomp_mutex_unlock (&devicep->lock);
+    return gomp_target_data_fallback ();
 
   struct target_mem_desc *tgt
     = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
@@ -1048,6 +1085,24 @@ GOMP_target_data (int device, const void *unused, size_t mapnum,
 }
 
 void
+GOMP_target_data_41 (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
+		     unsigned short *kinds)
+{
+  struct gomp_device_descr *devicep = resolve_device (device);
+
+  if (devicep == NULL
+      || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
+    return gomp_target_data_fallback ();
+
+  struct target_mem_desc *tgt
+    = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
+		     false);
+  struct gomp_task_icv *icv = gomp_icv (true);
+  tgt->prev = icv->target_data;
+  icv->target_data = tgt;
+}
+
+void
 GOMP_target_end_data (void)
 {
   struct gomp_task_icv *icv = gomp_icv (false);
@@ -1069,15 +1124,58 @@ GOMP_target_update (int device, const void *unused, size_t mapnum,
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
     return;
 
-  gomp_mutex_lock (&devicep->lock);
-  if (!devicep->is_initialized)
-    gomp_init_device (devicep);
-  gomp_mutex_unlock (&devicep->lock);
-
   gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false);
 }
 
 void
+GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
+			     size_t *sizes, unsigned short *kinds)
+{
+  struct gomp_device_descr *devicep = resolve_device (device);
+
+  if (devicep == NULL
+      || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
+    return;
+
+  /* Determine if this is an "omp target enter data".  */
+  const int typemask = 0xff;
+  bool is_enter_data = false;
+  size_t i;
+  for (i = 0; i < mapnum; i++)
+    {
+      unsigned char kind = kinds[i] & typemask;
+
+      if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+	continue;
+
+      if (kind == GOMP_MAP_ALLOC
+	  || kind == GOMP_MAP_TO
+	  || kind == GOMP_MAP_ALWAYS_TO)
+	{
+	  is_enter_data = true;
+	  break;
+	}
+
+      if (kind == GOMP_MAP_FROM
+	  || kind == GOMP_MAP_ALWAYS_FROM
+	  || kind == GOMP_MAP_DELETE
+	  || kind == GOMP_MAP_RELEASE)
+	break;
+
+      gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x", kind);
+    }
+
+  if (is_enter_data)
+    {
+      /* TODO  */
+    }
+  else
+    {
+      /* TODO  */
+    }
+}
+
+void
 GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
 {
   if (thread_limit)


  -- Ilya



More information about the Gcc-patches mailing list