On 2020/9/1 9:37 PM, Chung-Lin Tang wrote:
his 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.
Hi Jakub,
you haven't reviewed this 3rd part yet, but still updating with a rebased patch
here.
I've removed the above mentioned XFAILed testcase from the patch, since it
actually
belongs in the structure element mapping patches instead of here.
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.
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index da7ac037dcd..0cc3f4d406b 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1162,10 +1162,10 @@ struct gomp_device_descr
/* 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 *,
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
@@ -403,7 +403,8 @@ acc_map_data (void *h, void *d, size_t s)
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;
@@ -572,7 +573,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void
*kinds, int 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;
@@ -1202,7 +1204,7 @@ goacc_enter_data_internal (struct gomp_device_descr
*acc_dev, size_t mapnum,
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);
diff --git a/libgomp/target.c b/libgomp/target.c
index 1a8c67c2df5..61dab064fae 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -683,7 +683,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
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;
@@ -1212,15 +1212,16 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
/* 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:
@@ -1415,7 +1416,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
/* 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;
@@ -2475,6 +2476,19 @@ gomp_exit_data (struct gomp_device_descr *devicep,
size_t mapnum,
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;
@@ -2512,7 +2526,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t
mapnum,
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);
@@ -2621,6 +2637,14 @@ GOMP_target_enter_exit_data (int device, size_t mapnum,
void **hostaddrs,
&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);
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;
+}
+
