Hi Thomas,
On 11/29/19 3:40 PM, Thomas Schwinge wrote:
"[Patch][OpenMP/OpenACC/Fortran] Fix mapping of optional (present|absent)
arguments" reminded me that still this behavioral change has not been
split out, cited below, that you described as "trivial".
I've just filed <https://gcc.gnu.org/PR92726> "OpenACC: 'NULL'-in ->
no-op, and/or 'NULL'-out", so please reference that one in the ChangeLog
updates.
So, eventually that'll be more than just
'libgomp/oacc-mem.c:update_dev_host', but I understand we need that one
now, for Fortran optional arguments support. Any other changes can then
be handled later (once the OpenACC specification changes have been
completed).
For the changes at hand, they are called as library for
update_device(_async) and update_self(_async). And via the 'update'
directives via the 'maps' alway_pointer, (force_)to and (force_)from.
'insert_pointer' is called via 'GOACC_enter_exit_data for 'enter' if the
pointer can already be found; hence, I think the latter is untestable.
Side note: the [update_(device|self)]_async variants are mentioned in
the libgomp.texi but not really documented; they only appear at the end
of "OpenACC Profiling Interface" – last block at:
https://gcc.gnu.org/onlinedocs/libgomp/OpenACC-Profiling-Interface.html
I have the feeling the _async variants should be properly documented in
the manual; especially, the permitted values for the third argument.
Please also add a new test case 'libgomp.oacc-c-c++-common/null-1.c' with
a "Test 'NULL'-in -> no-op, and/or 'NULL'-out" header, executing things […]
How about the attached patch?
Tobias
2019-12-02 Tobias Burnus <tob...@codesourcery.com>
Kwok Cheung Yeung <k...@codesourcery.com>
libgomp/
* oacc-mem.c (update_dev_host, gomp_acc_insert_pointer): Just return
if input it a NULL pointer.
* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Remove; dependent on
diagnostic of NULL pointer.
* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/update-2.c: New; check whether
NULL is handled with update.
* testsuite/libgomp.oacc-fortran/update-2.f90: Ditto.
libgomp/libgomp.texi | 4 ++
libgomp/oacc-mem.c | 9 ++++
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c | 51 ----------------------
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c | 49 ---------------------
libgomp/testsuite/libgomp.oacc-c-c++-common/update-2.c | 33 ++++++++++++++
libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 | 40 +++++++++++++++++
6 files changed, 86 insertions(+), 100 deletions(-)
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 6db895f6272..4532e366f4b 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -2536,6 +2536,8 @@ This function updates the device copy from the previously mapped host memory.
The host memory is specified with the host address @var{a} and a length of
@var{len} bytes.
+If @var{a} is the @code{NULL} pointer, this is a no-op.
+
In Fortran, two (2) forms are supported. In the first form, @var{a} specifies
a contiguous array section. The second form @var{a} specifies a variable or
array element and @var{len} specifies the length in bytes.
@@ -2569,6 +2571,8 @@ This function updates the host copy from the previously mapped device memory.
The host memory is specified with the host address @var{a} and a length of
@var{len} bytes.
+If @var{a} is the @code{NULL} pointer, this is a no-op.
+
In Fortran, two (2) forms are supported. In the first form, @var{a} specifies
a contiguous array section. The second form @var{a} specifies a variable or
array element and @var{len} specifies the length in bytes.
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index aafe88d3a14..55c195bd819 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -829,6 +829,12 @@ update_dev_host (int is_dev, void *h, size_t s, int async)
if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
return;
+ /* Fortran optional arguments that are non-present result in a
+ NULL host address here. This can safely be ignored as it is
+ not possible to 'update' a non-present optional argument. */
+ if (h == NULL)
+ return;
+
acc_prof_info prof_info;
acc_api_info api_info;
bool profiling_p = GOACC_PROFILING_SETUP_P (thr, &prof_info, &api_info);
@@ -899,6 +905,9 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (*hostaddrs == NULL)
+ return;
+
if (acc_is_present (*hostaddrs, *sizes))
{
splay_tree_key n;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
deleted file mode 100644
index 5db29124e9e..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
+++ /dev/null
@@ -1,51 +0,0 @@
-/* Exercise acc_update_device with a NULL data address on nvidia targets. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d = acc_copyin (h, N);
- if (!d)
- abort ();
-
- for (i = 0; i < N; i++)
- {
- h[i] = 0xab;
- }
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_update_device (0, N);
-
- acc_copyout (h, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != 0xab)
- abort ();
- }
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[^\n\r]*,256\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c
deleted file mode 100644
index c2140429cb1..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c
+++ /dev/null
@@ -1,49 +0,0 @@
-/* Exercise acc_update_self with a NULL data mapping on nvidia targets. */
-
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
-#include <string.h>
-#include <stdlib.h>
-#include <openacc.h>
-
-int
-main (int argc, char **argv)
-{
- const int N = 256;
- int i;
- unsigned char *h;
- void *d;
-
- h = (unsigned char *) malloc (N);
-
- for (i = 0; i < N; i++)
- {
- h[i] = i;
- }
-
- d = acc_copyin (h, N);
- if (!d)
- abort ();
-
- memset (&h[0], 0, N);
-
- fprintf (stderr, "CheCKpOInT\n");
- acc_update_self (0, N);
-
- for (i = 0; i < N; i++)
- {
- if (h[i] != i)
- abort ();
- }
-
- acc_delete (h, N);
-
- free (h);
-
- return 0;
-}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[^\n\r]*,256\\\] is not mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-2.c
new file mode 100644
index 00000000000..7b6f8500b5e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-2.c
@@ -0,0 +1,33 @@
+/* Updating NULL pointer shall not fail at run time.
+ Cf. PR libgomp/92726 */
+
+#include <openacc.h>
+
+int
+main ()
+{
+ float *ptr = NULL;
+
+ #pragma acc enter data copyin(ptr)
+ #pragma acc enter data present_or_copyin(ptr)
+ #pragma acc enter data create(ptr)
+ #pragma acc enter data pcreate(ptr)
+
+ #pragma acc update self(ptr)
+ #pragma acc update host(ptr)
+ #pragma acc update device(ptr)
+
+ acc_update_device (ptr, 0);
+ acc_update_device (NULL, 0L);
+
+ acc_update_device_async (ptr, 0, acc_async_sync);
+ acc_update_device_async (NULL, 0L, acc_async_sync);
+
+ acc_update_self (ptr, 0);
+ acc_update_self (NULL, 0L);
+
+ acc_update_self_async (ptr, 0, acc_async_sync);
+ acc_update_self_async (NULL, 0L, acc_async_sync);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
new file mode 100644
index 00000000000..632ceb9ca7d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
@@ -0,0 +1,40 @@
+! Updating NULL pointer shall not fail at run time
+! Cf. PR libgomp/92726
+!
+
+program main
+ use openacc
+ use iso_c_binding
+ implicit none (type, external)
+
+ real, pointer :: ptr
+ real, pointer :: ptr2(:,:)
+
+ nullify(ptr)
+ nullify(ptr2)
+
+ !$acc enter data copyin(ptr, ptr2)
+ !$acc enter data present_or_copyin(ptr, ptr2)
+ !$acc enter data create(ptr, ptr2)
+ !$acc enter data pcreate(ptr, ptr2)
+
+ !$acc update self(ptr, ptr2)
+ !$acc update host(ptr, ptr2)
+ !$acc update device(ptr, ptr2)
+
+ call acc_update_device (ptr, 0_c_int32_t)
+ call acc_update_device (ptr2, 0_c_int64_t)
+ call acc_update_device (c_null_ptr, 0)
+
+ call acc_update_device_async (ptr, 0_c_int32_t, acc_async_sync)
+ call acc_update_device_async (ptr, 0_c_int64_t, acc_async_sync)
+ call acc_update_device_async (c_null_ptr, 0, acc_async_sync)
+
+ call acc_update_self (ptr, 0_c_int32_t)
+ call acc_update_self (ptr, 0_c_int64_t)
+ call acc_update_self (c_null_ptr, 0)
+
+ call acc_update_self_async (ptr, 0_c_int32_t, acc_async_sync)
+ call acc_update_self_async (ptr, 0_c_int64_t, acc_async_sync)
+ call acc_update_self_async (c_null_ptr, 0, acc_async_sync)
+end program main
