[1/5,OpenACC] Allow NULL as an argument to OpenACC 2.6 directives

Message ID f335efac-abfa-f43e-2afc-fd1931740362@codesourcery.com
State New
Headers show
Series
  • Add support for Fortran optional arguments in OpenACC
Related show

Commit Message

Kwok Cheung Yeung July 12, 2019, 11:35 a.m.
Fortran pass-by-reference optional arguments behave much like normal 
Fortran arguments when lowered to GENERIC/GIMPLE, except they can be 
null (representing a non-present argument).

Some parts of libgomp (those dealing with updating mappings) currently 
do not expect to take a null address and fail. These need to be changed 
to deal with the null appropriately, by turning the operation into a 
no-op (as you never need to update a non-present argument).

	libgomp/
	* oacc-mem.c (update_dev_host): Return early if the host address
	is NULL.
	(gomp_acc_insert_pointer): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Remove.
	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
---
  libgomp/oacc-mem.c                                 |  9 ++++
  .../testsuite/libgomp.oacc-c-c++-common/lib-43.c   | 51 
----------------------
  .../testsuite/libgomp.oacc-c-c++-common/lib-47.c   | 49 
---------------------
  3 files changed, 9 insertions(+), 100 deletions(-)
  delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
  delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c

-- 
2.8.1

Comments

Thomas Schwinge Nov. 29, 2019, 2:40 p.m. | #1
Hi Tobias!

Reviewing your
<http://mid.mail-archive.com/8be82276-81b1-817c-fcd2-51f24f5fe2d2@codesourcery.com>
"[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).

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
like 'acc_update_device (NULL, [...])' etc. for everything that calls
'update_dev_host': 'acc_update_device', 'acc_update_device_async',
'acc_update_self', 'acc_update_self_async'.  These functions are also
called for OpenACC 'update' directives
('libgomp/oacc-parallel.c:GOACC_update'), but I suppose it's not possible
to construct an OpenACC 'update' directive conveying a 'NULL' pointer,
that is, something that would result in 'hostaddrs[i] == NULL'?  Likewise
for Fortran, I suppose.

In 'libgomp/libgomp.texi' then add a note to 'acc_update_device' (and
other relevant functions) like this:

    @@ -2586,6 +2586,9 @@ 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.
    +
     [...]


As for 'libgomp/oacc-mem.c:gomp_acc_insert_pointer', that's only called
for OpenACC 'enter data' directives
('libgomp/oacc-parallel.c:GOACC_enter_exit_data'), and specifically only
for 'GOMP_MAP_POINTER', 'GOMP_MAP_TO_PSET'.  Is there a way to construct
a test case that will result in a 'NULL' pointer there, other than via
Fortran optional arguments?  If not, then that hunk should be removed
here, and move into/stay in the Fortran optional arguments patch that
you've posted.

(And that said, Julian's got a patch pending review that gets rid of
'gomp_acc_insert_pointer' and other such black magic, yay.)


Grüße
 Thomas


On 2019-07-12T12:35:05+0100, Kwok Cheung Yeung <kcy@codesourcery.com> wrote:
> Fortran pass-by-reference optional arguments behave much like normal 

> Fortran arguments when lowered to GENERIC/GIMPLE, except they can be 

> null (representing a non-present argument).

>

> Some parts of libgomp (those dealing with updating mappings) currently 

> do not expect to take a null address and fail. These need to be changed 

> to deal with the null appropriately, by turning the operation into a 

> no-op (as you never need to update a non-present argument).

>

> 	libgomp/

> 	* oacc-mem.c (update_dev_host): Return early if the host address

> 	is NULL.

> 	(gomp_acc_insert_pointer): Likewise.

> 	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Remove.

> 	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.

> ---

>   libgomp/oacc-mem.c                                 |  9 ++++

>   .../testsuite/libgomp.oacc-c-c++-common/lib-43.c   | 51 ----------------------

>   .../testsuite/libgomp.oacc-c-c++-common/lib-47.c   | 49 ---------------------

>   3 files changed, 9 insertions(+), 100 deletions(-)

>   delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c

>   delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c

>

> diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c

> index 2f27100..8cc5120 100644

> --- a/libgomp/oacc-mem.c

> +++ b/libgomp/oacc-mem.c

> @@ -831,6 +831,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);

> @@ -901,6 +907,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 5db2912..0000000

> --- 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 c214042..0000000

> --- 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 "" } */
Tobias Burnus Dec. 2, 2019, 4:59 p.m. | #2
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  <tobias@codesourcery.com>
	    Kwok Cheung Yeung <kcy@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

Patch

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2f27100..8cc5120 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -831,6 +831,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);
@@ -901,6 +907,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 5db2912..0000000
--- 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 c214042..0000000
--- 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 "" } */