[resending this 3rd patch since keep not seeing it on the list,
pardon if this gets duplicated]
This patch is the changes to libgomp and testcases.
There is now (again) a need to indicate OpenACC/OpenMP and
an 'enter data' style directive, so the associated changes to
'enum gomp_map_vars_kind'.
There is a slight change in the logic of gomp_attach_pointer
handling, because for OpenMP there might be a non-offloaded
data clause that attempts an attachment but silently continues
in case the pointer is not mapped.
Also in the testcases, an XFAILed testcase for structure element
mapping is added. OpenMP 5.0 specifies that a element of the same
structure variable are allocated/deallocated in a uniform fashion,
but this hasn't been implemented yet in this patch.
Thanks,
Chung-Lin
libgomp/
* libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag
usable.
* oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to
'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'.
(goacc_enter_datum): Likewise for call to gomp_map_vars_async.
(goacc_enter_data_internal): Likewise.
* target.c (gomp_map_vars_internal): Change checks of
GOMP_MAP_VARS_ENTER_DATA
to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases.
(gomp_exit_data): Add handling of GOMP_MAP_DETACH.
(GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH.
* testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
* testsuite/libgomp.c-c++-common/struct-elem-1.c: New xfailed testcase.
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index f9080e9f70f..3b53c08ba4f 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1145,18 +1145,18 @@ struct gomp_device_descr
/* This is mutable because of its mutable target_data member. */
acc_dispatch_t openacc;
};
/* Kind of the pragma, for which gomp_map_vars () is called. */
enum gomp_map_vars_kind
{
- GOMP_MAP_VARS_OPENACC,
- GOMP_MAP_VARS_TARGET,
- GOMP_MAP_VARS_DATA,
- GOMP_MAP_VARS_ENTER_DATA
+ GOMP_MAP_VARS_OPENACC = 1,
+ GOMP_MAP_VARS_TARGET = 2,
+ GOMP_MAP_VARS_DATA = 4,
+ GOMP_MAP_VARS_ENTER_DATA = 8
};
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
unsigned short *);
struct gomp_coalesce_buf;
extern void gomp_copy_host2dev (struct gomp_device_descr *,
struct goacc_asyncqueue *, void *, const void *,
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 65757ab2ffc..8dc521ac6d6 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -399,15 +399,16 @@ acc_map_data (void *h, void *d, size_t s)
(int)s);
}
gomp_mutex_unlock (&acc_dev->lock);
struct target_mem_desc *tgt
= gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
- &kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ &kinds, true,
+ GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
splay_tree_key n = tgt->list[0].key;
assert (n);
assert (n->refcount == 1);
assert (n->dynamic_refcount == 0);
/* Special reference counting behavior. */
@@ -568,15 +569,16 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void
*kinds, int async)
gomp_mutex_unlock (&acc_dev->lock);
goacc_aq aq = get_goacc_asyncqueue (async);
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ kinds, true,
+ GOMP_MAP_VARS_OPENACC |
GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
assert (n);
assert (n->refcount == 1);
assert (n->dynamic_refcount == 0);
n->dynamic_refcount++;
@@ -1198,15 +1200,15 @@ goacc_enter_data_internal (struct gomp_device_descr
*acc_dev, size_t mapnum,
member in the group has a NULL pointer (e.g. a non-present
optional parameter). */
gomp_mutex_unlock (&acc_dev->lock);
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
&sizes[i], &kinds[i], true,
- GOMP_MAP_VARS_ENTER_DATA);
+ GOMP_MAP_VARS_OPENACC |
GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
gomp_mutex_lock (&acc_dev->lock);
for (size_t j = 0; j < tgt->list_count; j++)
{
n = tgt->list[j].key;
diff --git a/libgomp/target.c b/libgomp/target.c
index 3e292eb8c62..ea6f29325b8 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -664,15 +664,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
const int rshift = short_mapkind ? 8 : 3;
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
- tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
+ tgt->refcount = (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) ? 0 : 1;
tgt->device_descr = devicep;
tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
if (mapnum == 0)
{
tgt->tgt_start = 0;
@@ -1093,23 +1093,24 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
tgt->list[i].length = n->host_end - n->host_start;
tgt->list[i].copy_from = false;
tgt->list[i].always_copy_from = false;
tgt->list[i].is_attach = true;
/* OpenACC 'attach'/'detach' doesn't affect
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
+
+ gomp_attach_pointer (devicep, aq, mem_map, n,
+ (uintptr_t) hostaddrs[i], sizes[i],
+ cbufp);
}
- else
+ else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("outer struct not mapped for attach");
}
- gomp_attach_pointer (devicep, aq, mem_map, n,
- (uintptr_t) hostaddrs[i], sizes[i],
- cbufp);
continue;
}
default:
break;
}
splay_tree_key k = &array->key;
k->host_start = (uintptr_t) hostaddrs[i];
@@ -1291,15 +1292,15 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
cbuf.buf = NULL;
cbufp = NULL;
}
/* If the variable from "omp target enter data" map-list was already mapped,
tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
gomp_exit_data. */
- if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+ if ((pragma_kind & GOMP_MAP_VARS_ENTER_DATA) && tgt->refcount == 0)
{
free (tgt);
tgt = NULL;
}
gomp_mutex_unlock (&devicep->lock);
return tgt;
@@ -2338,14 +2339,27 @@ gomp_exit_data (struct gomp_device_descr *devicep,
size_t mapnum,
gomp_mutex_lock (&devicep->lock);
if (devicep->state == GOMP_DEVICE_FINALIZED)
{
gomp_mutex_unlock (&devicep->lock);
return;
}
+ for (i = 0; i < mapnum; i++)
+ if ((kinds[i] & typemask) == GOMP_MAP_DETACH)
+ {
+ struct splay_tree_key_s cur_node;
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+
+ if (n)
+ gomp_detach_pointer (devicep, NULL, n, (uintptr_t) hostaddrs[i],
+ false, NULL);
+ }
+
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:
@@ -2375,15 +2389,17 @@ gomp_exit_data (struct gomp_device_descr *devicep,
size_t mapnum,
gomp_copy_dev2host (devicep, NULL, (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)
gomp_remove_var (devicep, k);
+ break;
+ case GOMP_MAP_DETACH:
break;
default:
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
kind);
}
}
@@ -2483,14 +2499,22 @@ GOMP_target_enter_exit_data (int device, size_t mapnum,
void **hostaddrs,
for (j = i + 1; j < mapnum; j++)
if (!GOMP_MAP_POINTER_P (get_kind (true, kinds, j) & 0xff))
break;
gomp_map_vars (devicep, j-i, &hostaddrs[i], NULL, &sizes[i],
&kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
i += j - i - 1;
}
+ else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH)
+ {
+ /* An attach operation must be processed together with the mapped
+ base-pointer list item. */
+ gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+ true, GOMP_MAP_VARS_ENTER_DATA);
+ i += 1;
+ }
else
gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
true, GOMP_MAP_VARS_ENTER_DATA);
else
gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds);
}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c
b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c
new file mode 100644
index 00000000000..b8012d6046e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c
@@ -0,0 +1,56 @@
+#include <stdlib.h>
+
+struct S
+{
+ int a, b;
+ int *ptr;
+ int c, d;
+};
+typedef struct S S;
+
+#define N 10
+int main (void)
+{
+ /* Test to see if pointer attachment works, for scalar pointers,
+ and pointer fields in structures. */
+
+ int *ptr = (int *) malloc (sizeof (int) * N);
+ int *orig_ptr = ptr;
+
+ #pragma omp target map (ptr, ptr[:N])
+ {
+ for (int i = 0; i < N; i++)
+ ptr[i] = N - i;
+ }
+
+ if (ptr != orig_ptr)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr[i] != N - i)
+ abort ();
+
+ S s = { 0 };
+ s.ptr = ptr;
+ #pragma omp target map (s, s.ptr[:N])
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i] = i;
+
+ s.a = 1;
+ s.b = 2;
+ }
+
+ if (s.ptr != ptr)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (s.ptr[i] != i)
+ abort ();
+
+ if (s.a != 1 || s.b != 2 || s.c != 0 || s.d != 0)
+ abort ();
+
+ return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
new file mode 100644
index 00000000000..bc7c38eae0a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
@@ -0,0 +1,32 @@
+/* { dg-xfail-run-if "TODO OpenMP 5.0 structure element mapping" { *-*-* } {
"*" } { "" } } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+ int a, b;
+};
+typedef struct S S;
+
+int main (void)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ S s;
+ #pragma omp target enter data map (alloc: s.a, s.b)
+ #pragma omp target exit data map (release: s.b)
+
+ /* OpenMP 5.0 structure element mapping rules describe that elements of same
+ structure variable should allocate/deallocate in a uniform fashion, so
+ "s.a" should be removed together by above 'exit data'. */
+ if (omp_target_is_present (&s.a, d))
+ abort ();
+
+ return 0;
+}
+
