@@ -107,6 +107,12 @@ enum gomp_map_kind
#define GOMP_MAP_POINTER_P(X) \
((X) == GOMP_MAP_POINTER)
+#define GOMP_MAP_ALWAYS_TO_P(X) \
+ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+
+#define GOMP_MAP_ALWAYS_FROM_P(X) \
+ (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
@@ -636,6 +636,15 @@ typedef struct splay_tree_node_s *splay_tree_node;
typedef struct splay_tree_s *splay_tree;
typedef struct splay_tree_key_s *splay_tree_key;
+struct target_var_desc {
+ /* Splay key. */
+ splay_tree_key key;
+ /* True if data should be copied from device to host at the end. */
+ bool copy_from;
+ /* True if data always should be copied from device to host at the end. */
+ bool always_copy_from;
+};
+
struct target_mem_desc {
/* Reference count. */
uintptr_t refcount;
@@ -655,9 +664,9 @@ struct target_mem_desc {
/* Corresponding target device descriptor. */
struct gomp_device_descr *device_descr;
- /* List of splay keys to remove (or decrease refcount)
+ /* List of target items to remove (or decrease refcount)
at the end of region. */
- splay_tree_key list[];
+ struct target_var_desc list[];
};
struct splay_tree_key_s {
@@ -673,8 +682,6 @@ struct splay_tree_key_s {
uintptr_t refcount;
/* Asynchronous reference count. */
uintptr_t async_refcount;
- /* True if data should be copied from device to host at the end. */
- bool copy_from;
};
#include "splay-tree.h"
@@ -651,7 +651,7 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
}
if (force_copyfrom)
- t->list[0]->copy_from = 1;
+ t->list[0].copy_from = 1;
gomp_mutex_unlock (&acc_dev->lock);
@@ -135,8 +135,8 @@ GOACC_parallel (int device, void (*fn) (void *),
devaddrs = gomp_alloca (sizeof (void *) * mapnum);
for (i = 0; i < mapnum; i++)
- devaddrs[i] = (void *) (tgt->list[i]->tgt->tgt_start
- + tgt->list[i]->tgt_offset);
+ devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset);
acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, sizes, kinds,
num_gangs, num_workers, vector_length, async,
@@ -161,6 +161,12 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn,
(void *) newn->host_start, (void *) newn->host_end,
(void *) oldn->host_start, (void *) oldn->host_end);
}
+
+ if (GOMP_MAP_ALWAYS_TO_P (kind))
+ devicep->host2dev_func (devicep->target_id,
+ (void *) (oldn->tgt->tgt_start + oldn->tgt_offset),
+ (void *) newn->host_start,
+ newn->host_end - newn->host_start);
oldn->refcount++;
}
@@ -260,7 +266,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
int kind = get_kind (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL)
{
- tgt->list[i] = NULL;
+ tgt->list[i].key = NULL;
continue;
}
cur_node.host_start = (uintptr_t) hostaddrs[i];
@@ -271,12 +277,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
if (n)
{
- tgt->list[i] = 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);
}
else
{
- tgt->list[i] = NULL;
+ tgt->list[i].key = NULL;
size_t align = (size_t) 1 << (kind >> rshift);
not_found_cnt++;
@@ -297,7 +306,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
break;
else
{
- tgt->list[j] = NULL;
+ tgt->list[j].key = NULL;
i++;
}
}
@@ -345,7 +354,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
size_t j;
for (i = 0; i < mapnum; i++)
- if (tgt->list[i] == NULL)
+ if (tgt->list[i].key == NULL)
{
int kind = get_kind (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL)
@@ -359,18 +368,23 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n)
{
- tgt->list[i] = 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);
}
else
{
size_t align = (size_t) 1 << (kind >> rshift);
- tgt->list[i] = k;
+ tgt->list[i].key = k;
tgt_size = (tgt_size + align - 1) & ~(align - 1);
k->tgt = tgt;
k->tgt_offset = tgt_size;
tgt_size += k->host_end - k->host_start;
- k->copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
+ 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);
k->refcount = 1;
k->async_refcount = 0;
tgt->refcount++;
@@ -388,6 +402,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
case GOMP_MAP_TOFROM:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_TOFROM:
/* FIXME: Perhaps add some smarts, like if copying
several adjacent fields from host to target, use some
host buffer to avoid sending each var individually. */
@@ -420,7 +436,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
break;
else
{
- tgt->list[j] = k;
+ tgt->list[j].key = k;
k->refcount++;
gomp_map_pointer (tgt,
(uintptr_t) *(void **) hostaddrs[j],
@@ -472,11 +488,11 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
{
for (i = 0; i < mapnum; i++)
{
- if (tgt->list[i] == NULL)
+ if (tgt->list[i].key == NULL)
cur_node.tgt_offset = (uintptr_t) NULL;
else
- cur_node.tgt_offset = tgt->list[i]->tgt->tgt_start
- + tgt->list[i]->tgt_offset;
+ cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset;
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -516,17 +532,17 @@ gomp_copy_from_async (struct target_mem_desc *tgt)
gomp_mutex_lock (&devicep->lock);
for (i = 0; i < tgt->list_count; i++)
- if (tgt->list[i] == NULL)
+ if (tgt->list[i].key == NULL)
;
- else if (tgt->list[i]->refcount > 1)
+ else if (tgt->list[i].key->refcount > 1)
{
- tgt->list[i]->refcount--;
- tgt->list[i]->async_refcount++;
+ tgt->list[i].key->refcount--;
+ tgt->list[i].key->async_refcount++;
}
else
{
- splay_tree_key k = tgt->list[i];
- if (k->copy_from)
+ splay_tree_key k = tgt->list[i].key;
+ if (tgt->list[i].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);
@@ -554,25 +570,33 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
size_t i;
for (i = 0; i < tgt->list_count; i++)
- if (tgt->list[i] == NULL)
- ;
- else if (tgt->list[i]->refcount > 1)
- tgt->list[i]->refcount--;
- else if (tgt->list[i]->async_refcount > 0)
- tgt->list[i]->async_refcount--;
- else
- {
- splay_tree_key k = tgt->list[i];
- if (k->copy_from && do_copyfrom)
- devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
- (void *) (k->tgt->tgt_start + k->tgt_offset),
- k->host_end - k->host_start);
- splay_tree_remove (&devicep->mem_map, k);
- if (k->tgt->refcount > 1)
- k->tgt->refcount--;
- else
- gomp_unmap_tgt (k->tgt);
- }
+ {
+ splay_tree_key k = tgt->list[i].key;
+ if (k == NULL)
+ continue;
+
+ bool do_unmap = false;
+ if (k->refcount > 1)
+ k->refcount--;
+ else if (k->async_refcount > 0)
+ k->async_refcount--;
+ else
+ do_unmap = true;
+
+ 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);
+ if (do_unmap)
+ {
+ splay_tree_remove (&devicep->mem_map, k);
+ if (k->tgt->refcount > 1)
+ k->tgt->refcount--;
+ else
+ gomp_unmap_tgt (k->tgt);
+ }
+ }
if (tgt->refcount > 1)
tgt->refcount--;
@@ -699,7 +723,6 @@ gomp_offload_image_to_device (struct gomp_device_descr *devicep,
k->tgt_offset = target_table[i].start;
k->refcount = 1;
k->async_refcount = 0;
- k->copy_from = false;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
@@ -725,7 +748,6 @@ gomp_offload_image_to_device (struct gomp_device_descr *devicep,
k->tgt_offset = target_var->start;
k->refcount = 1;
k->async_refcount = 0;
- k->copy_from = false;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
new file mode 100644
@@ -0,0 +1,51 @@
+/* { dg-require-effective-target offload_device } */
+
+#include <assert.h>
+
+int main ()
+{
+ int aa = 0, bb = 0, cc = 0, dd = 0;
+
+ #pragma omp target data map(tofrom: aa) map(to: bb) map(from: cc, dd)
+ {
+ int ok;
+ aa = bb = cc = 1;
+
+ /* Set dd on target to 0 for the further check. */
+ #pragma omp target map(always to: dd)
+ { dd; }
+
+ dd = 1;
+ #pragma omp target map(tofrom: aa) map(always to: bb) \
+ map(always from: cc) map(to: dd) map(from: ok)
+ {
+ /* bb is always to, aa and dd are not. */
+ ok = (aa == 0) && (bb == 1) && (dd == 0);
+ aa = bb = cc = dd = 2;
+ }
+
+ assert (ok);
+ assert (aa == 1);
+ assert (bb == 1);
+ assert (cc == 2); /* cc is always from. */
+ assert (dd == 1);
+
+ dd = 3;
+ #pragma omp target map(from: cc) map(always to: dd) map(from: ok)
+ {
+ ok = (dd == 3); /* dd is always to. */
+ cc = dd = 4;
+ }
+
+ assert (ok);
+ assert (cc == 2);
+ assert (dd == 3);
+ }
+
+ assert (aa == 2);
+ assert (bb == 1);
+ assert (cc == 4);
+ assert (dd == 4);
+
+ return 0;
+}