@@ -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 {
@@ -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);
@@ -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;
}