On Wed, Jul 29, 2015 at 10:06:52PM +0300, Ilya Verbin wrote:
> @@ -1241,6 +1245,62 @@ GOMP_target_update (int device, const void *unused,
> size_t mapnum,
> gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false);
> }
>
> +static void
> +gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
> + void **hostaddrs, size_t *sizes, unsigned short *kinds)
> +{
> + const int typemask = 0xff;
> + size_t i;
> + gomp_mutex_lock (&devicep->lock);
> + 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:
Please handle here GOMP_MAP_ZERO_LEN_ARRAY_SECTION too.
It should use gomp_map_lookup (while all others splay_tree_lookup),
otherwise it is the same as GOMP_MAP_RELEASE.
> @@ -1280,13 +1337,20 @@ GOMP_target_enter_exit_data (int device, size_t
> mapnum, void **hostaddrs,
> }
>
> if (is_enter_data)
> - {
> - /* TODO */
> - }
> + for (i = 0; i < mapnum; i++)
> + {
> + struct target_mem_desc *tgt_var
> + = gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i],
> + &kinds[i], true, false);
> + tgt_var->refcount--;
> +
> + /* If the variable was already mapped, tgt_var is not needed. Otherwise
> + tgt_var will be freed by gomp_unmap_vars or gomp_exit_data. */
> + if (tgt_var->refcount == 0)
> + free (tgt_var);
This is racy, you don't hold the device lock here anymore, so you shouldn't
decrease refcounts or test it etc.
I think better would be to change the bool is_target argument to
gomp_map_vars into an enum, and use 3 values there for now
- GOMP_VARS_MAP_TARGET, GOMP_VARS_MAP_DATA, GOMP_VARS_MAP_ENTER_DATA or so,
and for GOMP_VARS_MAP_ENTER_DATA perform the decrement of refcount and
freeing if it is zero (but then also better return NULL).
> diff --git a/libgomp/testsuite/libgomp.c/target-20.c
> b/libgomp/testsuite/libgomp.c/target-20.c
> new file mode 100644
> index 0000000..ec7e245
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-20.c
> @@ -0,0 +1,111 @@
> +/* { dg-require-effective-target offload_device } */
This test will fail on HSA, you don't assume just that it doesn't
fallback to host, but also non-shared address space.
I think it would be better to start with some check for non-shared address
space, like:
/* This test relies on non-shared address space. Punt otherwise. */
void ensure_nonshared_as (void)
{
int a = 8;
#pragma omp target map(to:a)
{
a++;
}
if (a == 8)
exit (0);
}
And generally, it is better to have most of the tests not relying on
offloading only or even non-shared address space, so that we also test
shared address space and host fallback. But a few tests won't hurt...
Jakub
