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