On Mon, Jul 06, 2015 at 22:42:10 +0200, Jakub Jelinek wrote:
> As has been clarified on omp-lang, we actually shouldn't be mapping or
> unmapping the pointer and/or reference, only the array slice itself, except
> in target construct (and even for that it is changing from mapping to
> private + pointer assignment).
I've updated this patch. make check-target-libgomp passed.
libgomp/
* target.c (gomp_map_vars_existing): Fix target address for 'always to'
array sections.
(gomp_unmap_vars): Decrement k->refcount when it is 1 and
k->async_refcount is 0.
(gomp_offload_image_to_device): Set tgt's refcount to infinity.
(gomp_exit_data): New static function.
(GOMP_target_enter_exit_data): Support mapping/unmapping.
* testsuite/libgomp.c/target-11.c: Extend for testing 'always to' array
sections.
* testsuite/libgomp.c/target-20.c: New test.
diff --git a/libgomp/target.c b/libgomp/target.c
index ef74d43..ad375c9 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -191,7 +191,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
splay_tree_key oldn,
if (GOMP_MAP_ALWAYS_TO_P (kind))
devicep->host2dev_func (devicep->target_id,
- (void *) (oldn->tgt->tgt_start + oldn->tgt_offset),
+ (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ + newn->host_start - oldn->host_start),
(void *) newn->host_start,
newn->host_end - newn->host_start);
if (oldn->refcount != REFCOUNT_INFINITY)
@@ -664,15 +665,18 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool
do_copyfrom)
continue;
bool do_unmap = false;
- if (k->refcount > 1)
+ if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+ k->refcount--;
+ else if (k->refcount == 1)
{
- if (k->refcount != REFCOUNT_INFINITY)
- k->refcount--;
+ if (k->async_refcount > 0)
+ k->async_refcount--;
+ else
+ {
+ k->refcount--;
+ do_unmap = true;
+ }
}
- 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)
@@ -798,7 +802,7 @@ gomp_offload_image_to_device (struct gomp_device_descr
*devicep,
/* Insert host-target address mapping into splay tree. */
struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
tgt->array = gomp_malloc ((num_funcs + num_vars) * sizeof (*tgt->array));
- tgt->refcount = 1;
+ tgt->refcount = REFCOUNT_INFINITY;
tgt->tgt_start = 0;
tgt->tgt_end = 0;
tgt->to_free = NULL;
@@ -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:
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizes[i];
+ splay_tree_key k = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (!k)
+ continue;
+
+ if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY)
+ k->refcount--;
+ if (kind == GOMP_MAP_DELETE && k->refcount != REFCOUNT_INFINITY)
+ k->refcount = 0;
+
+ if ((kind == GOMP_MAP_FROM && k->refcount == 0)
+ || kind == GOMP_MAP_ALWAYS_FROM)
+ devicep->dev2host_func (devicep->target_id,
+ (void *) cur_node.host_start,
+ (void *) (k->tgt->tgt_start + k->tgt_offset
+ + cur_node.host_start
+ - k->host_start),
+ cur_node.host_end - cur_node.host_start);
+ 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);
+ }
+
+ break;
+ default:
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
+ kind);
+ }
+ }
+
+ gomp_mutex_unlock (&devicep->lock);
+}
+
void
GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
size_t *sizes, unsigned short *kinds)
@@ -1259,9 +1319,6 @@ GOMP_target_enter_exit_data (int device, size_t mapnum,
void **hostaddrs,
{
unsigned char kind = kinds[i] & typemask;
- if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
- continue;
-
if (kind == GOMP_MAP_ALLOC
|| kind == GOMP_MAP_TO
|| kind == GOMP_MAP_ALWAYS_TO)
@@ -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);
+ }
else
- {
- /* TODO */
- }
+ gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds);
}
void
diff --git a/libgomp/testsuite/libgomp.c/target-11.c
b/libgomp/testsuite/libgomp.c/target-11.c
index b86097a..98882f0 100644
--- a/libgomp/testsuite/libgomp.c/target-11.c
+++ b/libgomp/testsuite/libgomp.c/target-11.c
@@ -9,6 +9,17 @@ void test_array_section (int *p)
{
#pragma omp target data map(alloc: p[0:N])
{
+ int ok = 1;
+ for (int i = 10; i < 10 + 4; i++)
+ p[i] = 997 * i;
+
+ #pragma omp target map(always to:p[10:4]) map(tofrom: ok)
+ for (int i = 10; i < 10 + 4; i++)
+ if (p[i] != 997 * i)
+ ok = 0;
+
+ assert (ok);
+
#pragma omp target map(always from:p[7:9])
for (int i = 0; i < N; i++)
p[i] = i;
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 } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 40
+
+int sum;
+int var1 = 1;
+int var2 = 2;
+
+#pragma omp declare target
+int D[N];
+#pragma omp end declare target
+
+void enter_data (int *X)
+{
+ #pragma omp target enter data map(to: var1, var2, X[:N]) map(alloc: sum)
+}
+
+void exit_data_0 (int *D)
+{
+ #pragma omp target exit data map(delete: D[:N])
+}
+
+void exit_data_1 ()
+{
+ #pragma omp target exit data map(from: var1)
+}
+
+void exit_data_2 (int *X)
+{
+ #pragma omp target exit data map(from: var2) map(release: X[:N], sum)
+}
+
+void test_nested ()
+{
+ int X = 0, Y = 0, Z = 0;
+
+ #pragma omp target data map(from: X, Y, Z)
+ {
+ #pragma omp target data map(from: X, Y, Z)
+ {
+ #pragma omp target map(from: X, Y, Z)
+ X = Y = Z = 1337;
+ assert (X == 0);
+ assert (Y == 0);
+ assert (Z == 0);
+
+ #pragma omp target exit data map(from: X) map(release: Y)
+ assert (X == 0);
+ assert (Y == 0);
+
+ #pragma omp target exit data map(release: Y) map(delete: Z)
+ assert (Y == 0);
+ assert (Z == 0);
+ }
+ assert (X == 1337);
+ assert (Y == 0);
+ assert (Z == 0);
+
+ #pragma omp target map(from: X)
+ X = 2448;
+ assert (X == 2448);
+ assert (Y == 0);
+ assert (Z == 0);
+
+ X = 4896;
+ }
+ assert (X == 4896);
+ assert (Y == 0);
+ assert (Z == 0);
+}
+
+int main ()
+{
+ int *X = malloc (N * sizeof (int));
+ int *Y = malloc (N * sizeof (int));
+ X[10] = 10;
+ Y[20] = 20;
+ enter_data (X);
+
+ exit_data_0 (D); /* This should have no effect on D. */
+
+ #pragma omp target map(alloc: var1, var2, X[:N]) map(to: Y[:N]) \
+ map(always from: sum)
+ {
+ var1 += X[10];
+ var2 += Y[20];
+ sum = var1 + var2;
+ D[sum]++;
+ }
+
+ assert (var1 == 1);
+ assert (var2 == 2);
+ assert (sum == 33);
+
+ exit_data_1 ();
+ assert (var1 == 11);
+ assert (var2 == 2);
+
+ exit_data_2 (X);
+ assert (var2 == 22);
+
+ free (X);
+ free (Y);
+
+ test_nested ();
+
+ return 0;
+}
-- Ilya
