On Tue, Jun 30, 2015 at 14:57:02 +0200, Jakub Jelinek wrote:
> On Tue, Jun 30, 2015 at 03:19:30PM +0300, Ilya Verbin wrote:
> > --- a/libgomp/target.c
> > +++ b/libgomp/target.c
> > @@ -580,10 +581,16 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool
> > do_copyfrom)
> > bool do_unmap = false;
> > if (k->refcount > 1)
> > k->refcount--;
> > - else if (k->async_refcount > 0)
> > - k->async_refcount--;
> > - else
> > - do_unmap = true;
> > + else if (k->refcount == 1)
> > + {
> > + if (k->async_refcount > 0)
> > + k->async_refcount--;
> > + else
> > + {
> > + k->refcount--;
> > + do_unmap = true;
> > + }
> > + }
>
> What is the rationale of this hunk change?
Without whis change, when k->refcount == 1, do_unmap is true, but refcount is
not decremented. So, if gomp_unmap_vars is called multiple times (now it's
possible for 4.1), refcount will remain 1, and it will try to unmap k at each
next call, that is wrong. That's why I decrement refcount to zero, and do
nothing when hit gomp_unmap_vars next time with k->refcount == 0.
> BTW, we'll likely need to treat also refcount == INT_MAX as special (never
> decrease it), because I believe declare target vars are supposed to have
> refcount of infinity rather than just 2GB-1.
I'll add special refcount for declare target vars.
> > @@ -1160,13 +1167,61 @@ GOMP_target_enter_exit_data (int device, size_t
> > mapnum, void **hostaddrs,
> > }
> >
> > if (is_enter_data)
> > - {
> > - /* TODO */
> > - }
> > + gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true,
> > false);
>
> This will leak the return value. Either we need to arrange not to allocate
> it for enter data, or we need to assign it to some variable and free
> immediately (we don't want to perform the release operations for it).
But we can't not allocate or free immediately it, since it's used later through
splay_tree_key_s::tgt, e.g. here:
if (is_target)
{
for (i = 0; i < mapnum; i++)
{
if (tgt->list[i].key == NULL)
cur_node.tgt_offset = (uintptr_t) NULL;
else
cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
+ tgt->list[i].key->tgt_offset;
My plan was to free tgt here:
+ if (k->refcount == 0)
+ {
+ splay_tree_remove (&devicep->mem_map, k);
+ if (k->tgt->refcount > 1)
+ k->tgt->refcount--;
+ else
+ gomp_unmap_tgt (k->tgt);
+ }
But now I understood that this will work only for simple cases like:
#pragma omp target enter data ...
...
#pragma omp target exit data ...
And will leak e.g. in:
#pragma omp target data ...
{
#pragma omp target enter data ...
}
> > else
> > - {
> > - /* TODO */
> > - }
> > + for (i = 0; i < mapnum; i++)
> > + {
> > + struct splay_tree_key_s cur_node;
> > + unsigned char kind = kinds[i] & typemask;
> > + switch (kind)
> > + {
> > + case GOMP_MAP_FROM:
> > + case GOMP_MAP_ALWAYS_FROM:
> > + case GOMP_MAP_DELETE:
> > + case GOMP_MAP_RELEASE:
> > + cur_node.host_start = (uintptr_t) hostaddrs[i];
> > + cur_node.host_end = cur_node.host_start + sizes[i];
> > + gomp_mutex_lock (&devicep->lock);
>
> I don't really like locking the mutex for each map clause in exit data
> separately. Perhaps just add a gomp_exit_data function similar to
> gomp_map_vars that will run this loop and be surrounded by the locking,
> or do it inline, but with the lock/unlock around the whole loop.
> exit data construct must have at least one map clause, so it doesn't make
> sense not to lock immediately.
I'll move locks outside of the loop.
> > + splay_tree_key k = splay_tree_lookup (&devicep->mem_map, &cur_node);
> > + if (!k)
> > + {
> > + gomp_mutex_unlock (&devicep->lock);
> > + continue;
> > + }
> > +
> > + if (k->refcount > 0)
> > + k->refcount--;
> > + if (kind == GOMP_MAP_DELETE)
> > + k->refcount = 0;
>
> See above, I believe delete should not delete refcount == INT_MAX
> mappings.
Will do that.
-- Ilya
