Hi,
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.
make check-target-libgomp and bootstrap passed, ok for gomp-4_1-branch?
gcc/
* builtin-types.def (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR): Remove.
(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_target1.
(BUILT_IN_GOMP_TARGET_DATA): Replace GOMP_target_data with
GOMP_target_data1.
(BUILT_IN_GOMP_TARGET_UPDATE): Replace GOMP_target_update with
GOMP_target_update1.
(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): Always use unsigned short for map kinds.
gcc/fortran/
* types.def (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR): Remove.
(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_target1, GOMP_target_data1,
GOMP_target_update1, GOMP_target_enter_exit_data.
* libgomp_g.h: Declare GOMP_target1, GOMP_target_data1,
GOMP_target_update1, 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_target1): New function.
(gomp_target_data_fallback): New static function.
(GOMP_target_data): Move host fallback to the new function.
(GOMP_target_data1): New function.
(GOMP_target_update): Do not call gomp_init_device.
(GOMP_target_update1): New function.
(GOMP_target_enter_exit_data): New function.
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 492ca63..3c4b9e3 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -524,8 +524,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_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 +535,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..18f81e6 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -189,8 +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_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,
@@ -199,9 +200,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..b8623af 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_target1",
+ BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data1",
+ 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_UPDATE, "GOMP_target_update1",
+ BT_FN_VOID_INT_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..3f39e1b 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,6 @@ 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 ();
- }
args.quick_push (t1);
args.quick_push (t2);
args.quick_push (t3);
@@ -10378,6 +10361,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:
{
@@ -12631,18 +12615,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
- tree tkind_type;
- int talign_shift;
- if (is_gimple_omp_oacc (stmt))
- {
- tkind_type = short_unsigned_type_node;
- talign_shift = 8;
- }
- else
- {
- tkind_type = unsigned_char_type_node;
- talign_shift = 3;
- }
+ tree tkind_type = short_unsigned_type_node;
+ int talign_shift = 8;
TREE_VEC_ELT (t, 2)
= create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
".omp_data_kinds");
@@ -12782,9 +12756,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..d15b9ba 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -242,6 +242,10 @@ GOMP_4.0.1 {
GOMP_4.1 {
global:
+ GOMP_target1;
+ GOMP_target_data1;
+ GOMP_target_update1;
+ 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..2536ee4 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -208,11 +208,19 @@ 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_target1 (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_data1 (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_update1 (int, size_t, void **, size_t *,
+ unsigned short *);
+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..bedc95a 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[1] 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[1] 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_target1 (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[1] 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_data1 (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,71 @@ 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_update1 (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_update (devicep, mapnum, hostaddrs, sizes, kinds, true);
+}
+
+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
