On Wed, Jun 24, 2015 at 13:39:03 +0200, Jakub Jelinek wrote:
> 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.
Indeed, here is the fix, make check-target-libgomp passed.
libgomp/
* libgomp.h (struct target_var_desc): Add offset and length.
* target.c (gomp_map_vars_existing): New argument tgt_var, fill it.
(gomp_map_vars): Move filling of tgt->list[i] into
gomp_map_vars_existing. Add missed case GOMP_MAP_ALWAYS_FROM.
(gomp_unmap_vars): Add list[i].offset to host and target addresses,
use list[i].length instead of k->host_end - k->host_start.
* testsuite/libgomp.c/target-11.c: Extend for testing array sections.
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index bd17828..c48e708 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -644,6 +644,12 @@ struct target_var_desc {
bool copy_from;
/* True if data always should be copied from device to host at the end. */
bool always_copy_from;
+ /* Used for unmapping of array sections, can be nonzero only when
+ always_copy_from is true. */
+ uintptr_t offset;
+ /* Used for unmapping of array sections, can be less than the size of the
+ whole object only when always_copy_from is true. */
+ uintptr_t length;
};
struct target_mem_desc {
diff --git a/libgomp/target.c b/libgomp/target.c
index b1640c1..a394e95 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -149,8 +149,15 @@ resolve_device (int device_id)
static inline void
gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn,
- splay_tree_key newn, unsigned char kind)
+ splay_tree_key newn, struct target_var_desc *tgt_var,
+ unsigned char kind)
{
+ tgt_var->key = oldn;
+ tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
+ tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
+ tgt_var->offset = newn->host_start - oldn->host_start;
+ tgt_var->length = newn->host_end - newn->host_start;
+
if ((kind & GOMP_MAP_FLAG_FORCE)
|| oldn->host_start > newn->host_start
|| oldn->host_end < newn->host_end)
@@ -276,13 +283,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
cur_node.host_end = cur_node.host_start + sizeof (void *);
splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
if (n)
- {
- tgt->list[i].key = n;
- tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
- tgt->list[i].always_copy_from
- = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
- gomp_map_vars_existing (devicep, n, &cur_node, kind & typemask);
- }
+ gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i],
+ kind & typemask);
else
{
tgt->list[i].key = NULL;
@@ -367,13 +369,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
k->host_end = k->host_start + sizeof (void *);
splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n)
- {
- tgt->list[i].key = n;
- tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
- tgt->list[i].always_copy_from
- = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
- gomp_map_vars_existing (devicep, n, k, kind & typemask);
- }
+ gomp_map_vars_existing (devicep, n, k, &tgt->list[i],
+ kind & typemask);
else
{
size_t align = (size_t) 1 << (kind >> rshift);
@@ -385,6 +382,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
tgt->list[i].always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+ tgt->list[i].offset = 0;
+ tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
k->async_refcount = 0;
tgt->refcount++;
@@ -397,6 +396,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
break;
case GOMP_MAP_TO:
case GOMP_MAP_TOFROM:
@@ -587,9 +587,11 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool
do_copyfrom)
if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
|| tgt->list[i].always_copy_from)
- devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
- (void *) (k->tgt->tgt_start + k->tgt_offset),
- k->host_end - k->host_start);
+ devicep->dev2host_func (devicep->target_id,
+ (void *) (k->host_start + tgt->list[i].offset),
+ (void *) (k->tgt->tgt_start + k->tgt_offset
+ + tgt->list[i].offset),
+ tgt->list[i].length);
if (do_unmap)
{
splay_tree_remove (&devicep->mem_map, k);
diff --git a/libgomp/testsuite/libgomp.c/target-11.c
b/libgomp/testsuite/libgomp.c/target-11.c
index 0fd183b..b86097a 100644
--- a/libgomp/testsuite/libgomp.c/target-11.c
+++ b/libgomp/testsuite/libgomp.c/target-11.c
@@ -1,7 +1,20 @@
/* { dg-require-effective-target offload_device } */
+#include <stdlib.h>
#include <assert.h>
+#define N 32
+
+void test_array_section (int *p)
+{
+ #pragma omp target data map(alloc: p[0:N])
+ {
+ #pragma omp target map(always from:p[7:9])
+ for (int i = 0; i < N; i++)
+ p[i] = i;
+ }
+}
+
int main ()
{
int aa = 0, bb = 0, cc = 0, dd = 0;
@@ -47,5 +60,16 @@ int main ()
assert (cc == 4);
assert (dd == 4);
+ int *array = calloc (N, sizeof (int));
+ test_array_section (array);
+
+ for (int i = 0; i < 7; i++)
+ assert (array[i] == 0);
+ for (int i = 7; i < 7 + 9; i++)
+ assert (array[i] == i);
+ for (int i = 7 + 9; i < N; i++)
+ assert (array[i] == 0);
+
+ free (array);
return 0;
}
-- Ilya
