From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 107553 invoked by alias); 15 Jun 2015 19:49:10 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 107541 invoked by uid 89); 15 Jun 2015 19:49:10 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=AWL,BAYES_00,FREEMAIL_FROM,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mail-qk0-f171.google.com Received: from mail-qk0-f171.google.com (HELO mail-qk0-f171.google.com) (209.85.220.171) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Mon, 15 Jun 2015 19:49:07 +0000 Received: by qkhu186 with SMTP id u186so11707764qkh.0 for ; Mon, 15 Jun 2015 12:49:04 -0700 (PDT) X-Received: by 10.140.46.75 with SMTP id j69mr37758480qga.17.1434397744742; Mon, 15 Jun 2015 12:49:04 -0700 (PDT) Received: from msticlxl57.ims.intel.com ([192.55.54.40]) by mx.google.com with ESMTPSA id f18sm6787625qka.9.2015.06.15.12.48.59 (version=TLSv1 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Mon, 15 Jun 2015 12:49:03 -0700 (PDT) Date: Mon, 15 Jun 2015 19:54:00 -0000 From: Ilya Verbin To: Jakub Jelinek Cc: Thomas Schwinge , gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data Message-ID: <20150615194850.GC45068@msticlxl57.ims.intel.com> References: <20150615122037.GA45068@msticlxl57.ims.intel.com> <20150615130609.GR10247@tucnak.redhat.com> <20150615161827.GB45068@msticlxl57.ims.intel.com> <20150615162528.GU10247@tucnak.redhat.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150615162528.GU10247@tucnak.redhat.com> User-Agent: Mutt/1.5.21 (2010-09-15) X-IsSubscribed: yes X-SW-Source: 2015-06/txt/msg01042.txt.bz2 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