OpenACC: Support GOMP_MAP_ZERO_LEN_ARRAY_SECTION
The shared code with OpenMP use special map kinds for zero-length arrays
(detected at runtime), but the OpenACC specific code doesn't know what to do
with them.
This patch implements support for GOMP_MAP_ZERO_LEN_ARRAY_SECTION and
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION throughout.
The last remaining problem case -- acc_is_present not reporting the array
present -- is highlighted in the testcase so it doesn't get forgotten, but will
need to be solved another time.
libgomp/ChangeLog:
* libgomp.h (splay_compare): Ensure that distinct zero-length mappings
aren't confused.
* oacc-mem.c (acc_is_present): Don't reject zero-sized queries.
(goacc_enter_datum): Likewise.
(update_dev_host): Don't actual copy zero-length arrays.
(goacc_enter_data_internal): Allow tgt to be null.
(goacc_exit_data_internal): Handle GOMP_MAP_ZERO_LEN_ARRAY_SECTION and
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION.
* oacc-parallel.c (GOACC_update): Handle
GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
* testsuite/libgomp.oacc-c/zerolengtharray.c: New test.
@@ -1026,7 +1026,8 @@ struct splay_tree_key_s {
static inline int
splay_compare (splay_tree_key x, splay_tree_key y)
{
- if (x->host_start == x->host_end
+ if (x->host_start == y->host_start
+ && x->host_start == x->host_end
&& y->host_start == y->host_end)
return 0;
if (x->host_end <= y->host_start)
@@ -322,7 +322,7 @@ acc_is_present (void *h, size_t s)
{
splay_tree_key n;
- if (!s || !h)
+ if (!h)
return 0;
goacc_lazy_initialize ();
@@ -534,7 +534,7 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
void *d;
splay_tree_key n;
- if (!hostaddrs[0] || !sizes[0])
+ if (!hostaddrs[0])
gomp_fatal ("[%p,+%d] is a bad range", hostaddrs[0], (int) sizes[0]);
goacc_lazy_initialize ();
@@ -849,6 +849,10 @@ update_dev_host (int is_dev, void *h, size_t s, int async)
if (h == NULL)
return;
+ /* Zero length arrays registered via gomp_map_vars don't show as mapped. */
+ if (s == 0)
+ return;
+
acc_prof_info prof_info;
acc_api_info api_info;
bool profiling_p = GOACC_PROFILING_SETUP_P (thr, &prof_info, &api_info);
@@ -1203,16 +1207,17 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
= gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
&sizes[i], &kinds[i], true,
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;
- if (n && !tgt->list[j].is_attach)
- n->dynamic_refcount++;
- }
+ /* TGT can be null for zero-length arrays. */
+ if (tgt)
+ for (size_t j = 0; j < tgt->list_count; j++)
+ {
+ n = tgt->list[j].key;
+ if (n && !tgt->list[j].is_attach)
+ n->dynamic_refcount++;
+ }
}
i = group_last;
@@ -1276,6 +1281,8 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
case GOMP_MAP_POINTER:
case GOMP_MAP_DELETE:
case GOMP_MAP_RELEASE:
+ case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+ case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
{
struct splay_tree_key_s cur_node;
size_t size;
@@ -647,6 +647,7 @@ GOACC_update (int flags_m, size_t mapnum,
{
case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
+ case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
break;
case GOMP_MAP_ALWAYS_POINTER:
new file mode 100644
@@ -0,0 +1,78 @@
+/* Ensure that GOMP_MAP_ZERO_LEN_ARRAY_SECTION is supported. */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+#ifndef ACC_MEM_SHARED
+#define ACC_MEM_SHARED 0
+#endif
+
+int a[100];
+int b[100];
+int c[100];
+int d[100];
+int e[100];
+
+int s = 0; // runtime size is zero
+
+int main ()
+{
+ /* Check it's not present too early.
+ Host fallback and shared-memory systems will show the data present. */
+ if (acc_is_present (a, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (acc_is_present (b, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (acc_is_present (d, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (acc_is_present (e, s) == !ACC_MEM_SHARED)
+ abort ();
+
+ /* Test GOMP_MAP_ZERO_LEN_ARRAY_SECTION inputs. */
+#pragma acc enter data create(a[0:s])
+#pragma acc enter data copyin(b[0:s])
+ acc_create (d, s);
+ acc_copyin (e, s);
+#pragma acc update device(a[0:s])
+ acc_update_device (b, s);
+
+ // FIXME: update these when they work correctly
+ // TODO { dg-output "acc_is_present does not currently work for zero-length arrays created via pragmas" { xfail *-*-* } }
+ if (/*!*/acc_is_present (a, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (/*!*/acc_is_present (b, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (!acc_is_present (d, s))
+ abort ();
+ if (!acc_is_present (e, s))
+ abort ();
+
+#pragma acc parallel copy(c[0:s])
+ {
+ ;
+ }
+
+ /* Test GOMP_MAP_ZERO_LEN_ARRAY_SECTION outputs. */
+ acc_update_self (a, s);
+#pragma acc update self(b[0:s])
+ acc_copyout (e, s);
+#pragma acc exit data copyout(b[0:s])
+
+ /* Test GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION. */
+ acc_delete (d, s);
+#pragma acc exit data delete(a[0:s]) finalize
+
+ /* Check it's no longer present. */
+ if (acc_is_present (a, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (acc_is_present (b, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (acc_is_present (d, s) == !ACC_MEM_SHARED)
+ abort ();
+ if (acc_is_present (e, s) == !ACC_MEM_SHARED)
+ abort ();
+
+ return 0;
+}