[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