This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Ilya Verbin <iverbin at gmail dot com>
- Cc: Thomas Schwinge <thomas at codesourcery dot com>, gcc-patches at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>
- Date: Wed, 24 Jun 2015 13:39:03 +0200
- Subject: Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Authentication-results: sourceware.org; auth=none
- References: <20150615122037 dot GA45068 at msticlxl57 dot ims dot intel dot com> <20150615130609 dot GR10247 at tucnak dot redhat dot com> <20150615161827 dot GB45068 at msticlxl57 dot ims dot intel dot com> <20150615162528 dot GU10247 at tucnak dot redhat dot com> <20150615194850 dot GC45068 at msticlxl57 dot ims dot intel dot com> <20150615195840 dot GZ10247 at tucnak dot redhat dot com> <20150619213514 dot GA23723 at msticlxl57 dot ims dot intel dot com> <20150623114043 dot GC18789 at msticlxl57 dot ims dot intel dot com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Tue, Jun 23, 2015 at 02:40:43PM +0300, Ilya Verbin wrote:
> On Sat, Jun 20, 2015 at 00:35:14 +0300, Ilya Verbin wrote:
> > Given that a mapped variable in 4.1 can have different kinds across nested data
> > regions, we need to store map-type not only for each var, but also for each
> > structured mapping. Here is my WIP patch, is it sane? :)
> > Attached testcase works OK on the device with non-shared memory.
>
> A bit updated version with a fix for GOMP_MAP_TO_PSET.
> make check-target-libgomp passed.
Thinking about this more, for always modifier this isn't really sufficient.
Consider:
void
foo (int *p)
{
#pragma omp target data (alloc:p[0:32])
{
#pragma omp target data (always, from:p[7:9])
{
...
}
}
}
If all we record is the corresponding splay_tree and the flags
(from/always_from), then this would try to copy from the device
the whole array section, rather than just the small portion of it.
So, supposedly in addition to the splay_tree for always from case we also
need to remember e.g. [relative offset, length] within the splay tree
object.
Jakub
- References:
- [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data
- Re: [gomp4.1] Add new versions of GOMP_target{,_data,_update} and GOMP_target_enter_exit_data