From 84217e9d4a08d206880f9dad6a8065bf2bc00a2d Mon Sep 17 00:00:00 2001 From: jakub Date: Tue, 12 Apr 2016 07:05:29 +0000 Subject: * omp-low.c (lower_omp_target): Use GOMP_MAP_FIRSTPRIVATE_INT regardless whether there are depend clauses or not. * libgomp.h (struct gomp_target_task): Remove firstprivate_copies field. * target.c (gomp_target_fallback_firstprivate, gomp_target_unshare_firstprivate): Removed. (GOMP_target_ext): Copy firstprivate vars into gomp_allocaed memory before waiting for dependencies. (gomp_target_task_fn): Don't copy firstprivate vars here. * task.c (GOMP_PLUGIN_target_task_completion): Don't free firstprivate_copies here. (gomp_create_target_task): Don't initialize firstprivate_copies field. * testsuite/libgomp.c/target-25.c (main): Use map (to:) instead of explicit/implicit firstprivate. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@234894 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 15 ++++++ libgomp/libgomp.h | 2 - libgomp/target.c | 92 ++++++++++++++------------------- libgomp/task.c | 2 - libgomp/testsuite/libgomp.c/target-25.c | 2 +- 5 files changed, 55 insertions(+), 58 deletions(-) (limited to 'libgomp') diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 742f19052a4..b53dc6b7056 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,18 @@ +2016-04-12 Jakub Jelinek + + * libgomp.h (struct gomp_target_task): Remove firstprivate_copies + field. + * target.c (gomp_target_fallback_firstprivate, + gomp_target_unshare_firstprivate): Removed. + (GOMP_target_ext): Copy firstprivate vars into gomp_allocaed memory + before waiting for dependencies. + (gomp_target_task_fn): Don't copy firstprivate vars here. + * task.c (GOMP_PLUGIN_target_task_completion): Don't free + firstprivate_copies here. + (gomp_create_target_task): Don't initialize firstprivate_copies field. + * testsuite/libgomp.c/target-25.c (main): Use map (to:) instead of + explicit/implicit firstprivate. + 2016-04-08 Cesar Philippidis PR lto/70289 diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 7108a6d0118..664e76b52d1 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -496,8 +496,6 @@ struct gomp_target_task struct target_mem_desc *tgt; struct gomp_task *task; struct gomp_team *team; - /* Copies of firstprivate mapped data for shared memory accelerators. */ - void *firstprivate_copies; /* Device-specific target arguments. */ void **args; void *hostaddrs[]; diff --git a/libgomp/target.c b/libgomp/target.c index 96fe3d5eb0d..e2dd0e08997 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1372,47 +1372,6 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs, } } -/* Host fallback with firstprivate map-type handling. */ - -static void -gomp_target_fallback_firstprivate (void (*fn) (void *), size_t mapnum, - void **hostaddrs, size_t *sizes, - unsigned short *kinds) -{ - size_t tgt_align = 0, tgt_size = 0; - calculate_firstprivate_requirements (mapnum, sizes, kinds, &tgt_align, - &tgt_size); - if (tgt_align) - { - char *tgt = gomp_alloca (tgt_size + tgt_align - 1); - copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, tgt_align, - tgt_size); - } - gomp_target_fallback (fn, hostaddrs); -} - -/* Handle firstprivate map-type for shared memory devices and the host - fallback. Return the pointer of firstprivate copies which has to be freed - after use. */ - -static void * -gomp_target_unshare_firstprivate (size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds) -{ - size_t tgt_align = 0, tgt_size = 0; - char *tgt = NULL; - - calculate_firstprivate_requirements (mapnum, sizes, kinds, &tgt_align, - &tgt_size); - if (tgt_align) - { - tgt = gomp_malloc (tgt_size + tgt_align - 1); - copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, tgt_align, - tgt_size); - } - return tgt; -} - /* Helper function of GOMP_target{,_ext} routines. */ static void * @@ -1504,6 +1463,8 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, unsigned int flags, void **depend, void **args) { struct gomp_device_descr *devicep = resolve_device (device); + size_t tgt_align = 0, tgt_size = 0; + bool fpc_done = false; if (flags & GOMP_TARGET_FLAG_NOWAIT) { @@ -1555,7 +1516,19 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, { struct gomp_thread *thr = gomp_thread (); if (thr->task && thr->task->depend_hash) - gomp_task_maybe_wait_for_dependencies (depend); + { + /* If we might need to wait, copy firstprivate now. */ + calculate_firstprivate_requirements (mapnum, sizes, kinds, + &tgt_align, &tgt_size); + if (tgt_align) + { + char *tgt = gomp_alloca (tgt_size + tgt_align - 1); + copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, + tgt_align, tgt_size); + } + fpc_done = true; + gomp_task_maybe_wait_for_dependencies (depend); + } } void *fn_addr; @@ -1564,15 +1537,35 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, || !(fn_addr = gomp_get_target_fn_addr (devicep, fn)) || (devicep->can_run_func && !devicep->can_run_func (fn_addr))) { - gomp_target_fallback_firstprivate (fn, mapnum, hostaddrs, sizes, kinds); + if (!fpc_done) + { + calculate_firstprivate_requirements (mapnum, sizes, kinds, + &tgt_align, &tgt_size); + if (tgt_align) + { + char *tgt = gomp_alloca (tgt_size + tgt_align - 1); + copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, + tgt_align, tgt_size); + } + } + gomp_target_fallback (fn, hostaddrs); return; } struct target_mem_desc *tgt_vars; - void *fpc = NULL; if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) { - fpc = gomp_target_unshare_firstprivate (mapnum, hostaddrs, sizes, kinds); + if (!fpc_done) + { + calculate_firstprivate_requirements (mapnum, sizes, kinds, + &tgt_align, &tgt_size); + if (tgt_align) + { + char *tgt = gomp_alloca (tgt_size + tgt_align - 1); + copy_firstprivate_data (tgt, mapnum, hostaddrs, sizes, kinds, + tgt_align, tgt_size); + } + } tgt_vars = NULL; } else @@ -1583,8 +1576,6 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, args); if (tgt_vars) gomp_unmap_vars (tgt_vars, true); - else - free (fpc); } /* Host fallback for GOMP_target_data{,_ext} routines. */ @@ -1891,9 +1882,7 @@ gomp_target_task_fn (void *data) || (devicep->can_run_func && !devicep->can_run_func (fn_addr))) { ttask->state = GOMP_TARGET_TASK_FALLBACK; - gomp_target_fallback_firstprivate (ttask->fn, ttask->mapnum, - ttask->hostaddrs, ttask->sizes, - ttask->kinds); + gomp_target_fallback (ttask->fn, ttask->hostaddrs); return false; } @@ -1908,9 +1897,6 @@ gomp_target_task_fn (void *data) if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) { ttask->tgt = NULL; - ttask->firstprivate_copies - = gomp_target_unshare_firstprivate (ttask->mapnum, ttask->hostaddrs, - ttask->sizes, ttask->kinds); actual_arguments = ttask->hostaddrs; } else diff --git a/libgomp/task.c b/libgomp/task.c index 38d4e9b413b..023663f43d5 100644 --- a/libgomp/task.c +++ b/libgomp/task.c @@ -582,7 +582,6 @@ GOMP_PLUGIN_target_task_completion (void *data) return; } ttask->state = GOMP_TARGET_TASK_FINISHED; - free (ttask->firstprivate_copies); gomp_target_task_completion (team, task); gomp_mutex_unlock (&team->task_lock); } @@ -683,7 +682,6 @@ gomp_create_target_task (struct gomp_device_descr *devicep, ttask->state = state; ttask->task = task; ttask->team = team; - ttask->firstprivate_copies = NULL; task->fn = NULL; task->fn_data = ttask; task->final_task = 0; diff --git a/libgomp/testsuite/libgomp.c/target-25.c b/libgomp/testsuite/libgomp.c/target-25.c index aeb19aee510..09b8d52184a 100644 --- a/libgomp/testsuite/libgomp.c/target-25.c +++ b/libgomp/testsuite/libgomp.c/target-25.c @@ -23,7 +23,7 @@ main () usleep (7000); z = 3; } - #pragma omp target map(tofrom: x) map(from: err) firstprivate (y) depend(inout: x, z) + #pragma omp target map(tofrom: x) map(from: err) map (to: y, z) depend(inout: x, z) err = (x != 1 || y != 2 || z != 3); if (err) abort (); -- cgit v1.2.1