diff mbox series

OpenACC: Support GOMP_MAP_ZERO_LEN_ARRAY_SECTION

Message ID cb91db36-8762-c798-2ccb-411b743de1ef@codesourcery.com
State New
Headers show
Series OpenACC: Support GOMP_MAP_ZERO_LEN_ARRAY_SECTION | expand

Commit Message

Andrew Stubbs July 29, 2020, 2:14 p.m. UTC 
This patch adds support for zero-length arrays in OpenACC data 
transfers. Previously, trying to use an array section with zero length 
would cause a fatal error at runtime.

This patch requires that my other patch "OpenACC: Separate enter/exit 
data APIs" is already applied.

Unfortunately, because the reference counting is handled by the code 
shared with OpenMP, and because the semantics there appear to be a 
little bit different (or broken?), I've been unable to get 
acc_is_present to return true for zero-length arrays created by pragmas 
(those created via acc_create are fine). That issue will require a 
another patch, probably with more invasive changes.

The test case should cover all the main uses of zero-length arrays, and 
I've added an xfail message to highlight the known deficiency.

OK for mainline (and backport to OG10)?

Andrew
diff mbox series

Patch

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.

diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index f9080e9f70f..e0426acdbfe 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -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)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 45162d24786..965c81ddbd7 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -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;
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index bca31b51427..d3277e60404 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -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:
diff --git a/libgomp/testsuite/libgomp.oacc-c/zerolengtharray.c b/libgomp/testsuite/libgomp.oacc-c/zerolengtharray.c
new file mode 100644
index 00000000000..cae102cb580
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c/zerolengtharray.c
@@ -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;
+}