| Message ID | c6b6ab11-0113-a51e-d8c6-ea4d83332f7f@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | None | expand |
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; +} +
On 2020/10/28 6:33 PM, Chung-Lin Tang wrote: > 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. For the libgomp patch, v3 doesn't update any of the code proper, but the libgomp.c-c++-common/ptr-attach-1.c testcase had some code added to test the case of a base-pointer on device by "declare target". Thanks, Chung-Lin 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..e7deec6e006 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c @@ -0,0 +1,82 @@ +#include <stdlib.h> + +struct S +{ + int a, b; + int *ptr; + int c, d; +}; +typedef struct S S; + +#pragma omp declare target +int *gp; +#pragma omp end declare target + +#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 (); + + gp = (int *) malloc (sizeof (int) * N); + orig_ptr = gp; + + for (int i = 0; i < N; i++) + gp[i] = i - 1; + + #pragma omp target map (gp[:N]) + { + for (int i = 0; i < N; i++) + gp[i] += 1; + } + + if (gp != orig_ptr) + abort (); + + for (int i = 0; i < N; i++) + if (gp[i] != i) + abort (); + + free (ptr); + free (gp); + + return 0; +} +
On Wed, Nov 04, 2020 at 02:03:27AM +0800, Chung-Lin Tang wrote: > > Â Â Â Â 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. Ok, with two nits fixed. > @@ -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); This line is too long. > 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); And this one too. Please wrap them. Jakub
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; +} +
[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.