C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct.

This patch adds the 'has_device_addr' clause to the OpenMP 'target' construct
which was introduced in OpenMP 5.1 (OpenMP API 5.1 specification pp. 197ff):

	has_device_addr(list)

"The has_device_addr clause indicates that its list items already have device
addresses and therefore they may be directly accessed from a target device.
If the device address of a list item is not for the device on which the target
region executes, accessing the list item inside the region results in
unspecified behavior. The list items may include array sections." (p. 200)

"A list item may not be specified in both an is_device_ptr clause and a
has_device_addr clause on the directive." (p. 202)

"A list item that appears in an is_device_ptr or a has_device_addr clause must
not be specified in any data-sharing attribute clause on the same target
construct." (p. 203)

gcc/c-family/ChangeLog:

	* c-omp.cc (c_omp_split_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR case.
	* c-pragma.h (enum pragma_kind): Added 5.1 in comment.
	(enum pragma_omp_clause): Added PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_clause_name): Parse 'has_device_addr'
	clause.
	(c_parser_omp_variable_list): Handle array sections.
	(c_parser_omp_clause_has_device_addr): Added.
	(c_parser_omp_all_clauses): Added PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR
	case.
	(c_parser_omp_target_exit_data): Added HAS_DEVICE_ADDR to
	OMP_CLAUSE_MASK.
	* c-typeck.cc (handle_omp_array_sections): Handle clause restrictions.
	(c_finish_omp_clauses): Handle array sections.

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_omp_clause_name): Parse 'has_device_addr' clause.
	(cp_parser_omp_var_list_no_open): Handle array sections.
	(cp_parser_omp_all_clauses): Added PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR
	case.
	(cp_parser_omp_target_update): Added HAS_DEVICE_ADDR to OMP_CLAUSE_MASK.
	* semantics.cc (handle_omp_array_sections): Handle clause restrictions.
	(finish_omp_clauses): Handle array sections.

gcc/fortran/ChangeLog:

	* dump-parse-tree.cc (show_omp_clauses): Added OMP_LIST_HAS_DEVICE_ADDR
	case.
	* gfortran.h: Added OMP_LIST_HAS_DEVICE_ADDR.
	* openmp.cc (enum omp_mask2): Added OMP_CLAUSE_HAS_DEVICE_ADDR.
	(gfc_match_omp_clauses): Parse HAS_DEVICE_ADDR clause.
	(resolve_omp_clauses): Same.
	* trans-openmp.cc (gfc_trans_omp_variable_list): Added
	OMP_LIST_HAS_DEVICE_ADDR case.
	(gfc_trans_omp_clauses): Firstprivatize of array descriptors.

gcc/ChangeLog:

	* gimplify.cc (gimplify_scan_omp_clauses): Added cases for
	OMP_CLAUSE_HAS_DEVICE_ADDR
	and handle array sections.
	(gimplify_adjust_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR case.
	* omp-low.cc (scan_sharing_clauses): Handle OMP_CLAUSE_HAS_DEVICE_ADDR.
	(lower_omp_target): Same.
	* tree-core.h (enum omp_clause_code): Same.
	* tree-nested.cc (convert_nonlocal_omp_clauses): Same.
	(convert_local_omp_clauses): Same.
	* tree-pretty-print.cc (dump_omp_clause): Same.
	* tree.cc: Same.

libgomp/ChangeLog:

	* libgomp.texi: Updated entry for HAS_DEVICE_ADDR.
	* target.c (copy_firstprivate_data): Copy only if host address is not
	NULL.
	* testsuite/libgomp.c++/target-has-device-addr-2.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-4.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-5.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-6.C: New test.
	* testsuite/libgomp.c-c++-common/target-has-device-addr-1.c: New test.
	* testsuite/libgomp.c/target-has-device-addr-3.c: New test.
	* testsuite/libgomp.fortran/target-has-device-addr-1.f90: New test.
	* testsuite/libgomp.fortran/target-has-device-addr-2.f90: New test.
	* testsuite/libgomp.fortran/target-has-device-addr-3.f90: New test.
	* testsuite/libgomp.fortran/target-has-device-addr-4.f90: New test.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/clauses-1.c: Added has_device_addr to test cases.
	* g++.dg/gomp/attrs-1.C: Added has_device_addr to test cases.
	* g++.dg/gomp/attrs-2.C: Added has_device_addr to test cases.
	* c-c++-common/gomp/target-has-device-addr-1.c: New test.
	* c-c++-common/gomp/target-has-device-addr-2.c: New test.
	* c-c++-common/gomp/target-is-device-ptr-1.c: New test.
	* c-c++-common/gomp/target-is-device-ptr-2.c: New test.
	* gfortran.dg/gomp/is_device_ptr-3.f90: New test.
	* gfortran.dg/gomp/target-has-device-addr-1.f90: New test.
	* gfortran.dg/gomp/target-has-device-addr-2.f90: New test.

12 files changed, 480 insertions, 2 deletions

diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 424459f4442..161a423ac7c 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -294,7 +294,7 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{align} clause/modifier in @code{allocate} directive/clause
       and @code{allocator} directive @tab P @tab C/C++ on clause only
 @item @code{thread_limit} clause to @code{target} construct @tab Y @tab
-@item @code{has_device_addr} clause to @code{target} construct @tab N @tab
+@item @code{has_device_addr} clause to @code{target} construct @tab Y @tab
 @item iterators in @code{target update} motion clauses and @code{map}
       clauses @tab N @tab
 @item indirect calls to the device version of a procedure or function in
diff --git a/libgomp/target.c b/libgomp/target.c
index 698ff14a05f..9017458885e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2510,7 +2510,7 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs,
   tgt_size = 0;
   size_t i;
   for (i = 0; i < mapnum; i++)
-    if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+    if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE && hostaddrs[i] != NULL)
       {
 	size_t align = (size_t) 1 << (kinds[i] >> 8);
 	tgt_size = (tgt_size + align - 1) & ~(align - 1);
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C
new file mode 100644
index 00000000000..d9a309d7af4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C
@@ -0,0 +1,23 @@
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc (sizeof(int), 0);
+
+  #pragma omp target is_device_ptr(dp)
+    *dp = 42;
+
+  int &x = *dp;
+
+  #pragma omp target has_device_addr(x)
+    x = 24;
+
+  #pragma omp target has_device_addr(x)
+    if (x != 24)
+      __builtin_abort ();
+
+  omp_target_free(dp, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C
new file mode 100644
index 00000000000..6468c6c8433
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C
@@ -0,0 +1,33 @@
+#include <omp.h>
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc (30*sizeof(int), 0);
+
+  #pragma omp target is_device_ptr(dp)
+    for (int i = 0; i < 30; i++)
+      dp[i] = i;
+
+  int (&x)[30] = *static_cast<int(*)[30]>(static_cast<void*>(dp));
+
+  #pragma omp target has_device_addr(x)
+    for (int i = 0; i < 30; i++)
+      x[i] = 2 * i;
+
+  #pragma omp target has_device_addr(x)
+    for (int i = 0; i < 30; i++)
+      if (x[i] != 2 * i)
+	__builtin_abort ();
+
+  #pragma omp target has_device_addr(x[1:5])
+    for (int i = 1; i < 6; i++)
+      x[i] = 3 * i;
+
+  #pragma omp target has_device_addr(x[1:5])
+    for (int i = 1; i < 6; i++)
+      if (x[i] != 3 * i)
+	__builtin_abort ();
+
+  omp_target_free (dp, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C
new file mode 100644
index 00000000000..e847cdceb44
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C
@@ -0,0 +1,33 @@
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+  int *dpx = (int*)omp_target_alloc (sizeof(int), 0);
+  int **dpy = (int**)omp_target_alloc (sizeof(int*), 0);
+
+  #pragma omp target is_device_ptr(dpx, dpy)
+    {
+      *dpx = 42;
+      int z = 77;
+      *dpy = &z;
+    }
+
+  int& x = *dpx;
+  int*& y = *dpy;
+
+  #pragma omp target has_device_addr(x, y)
+    {
+      x = 24;
+      y = &x;
+    }
+
+  #pragma omp target has_device_addr(x, y)
+    if (x != 24 || y != &x)
+      __builtin_abort ();
+
+  omp_target_free(dpx, 0);
+  omp_target_free(dpy, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C
new file mode 100644
index 00000000000..141edb14dec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C
@@ -0,0 +1,32 @@
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+  int *dpx = (int*)omp_target_alloc (sizeof(int), 0);
+  double *dpy = (double*)omp_target_alloc (sizeof(double), 0);
+
+  #pragma omp target is_device_ptr(dpx, dpy)
+    {
+      *dpx = 42;
+      *dpy = 43.5;
+    }
+
+  int &x = *dpx;
+  double &y = *dpy;
+
+  #pragma omp target has_device_addr(x, y)
+    {
+      x = 24;
+      y = 25.7;
+    }
+
+  #pragma omp target has_device_addr(y, x)
+    if (x != 24 || y != 25.7)
+      __builtin_abort ();
+
+  omp_target_free(dpx, 0);
+  omp_target_free(dpy, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c
new file mode 100644
index 00000000000..fcc5c9e8553
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c
@@ -0,0 +1,73 @@
+/* Testing the 'has_device_addr' clause on the target construct with
+   enclosing 'target data' construct. */
+
+#define N 40
+
+int
+main ()
+{
+  int x = 24;
+
+  #pragma omp target data map(x) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      x = 42;
+  if (x != 42)
+    __builtin_abort ();
+
+  int y[N];
+
+  for (int i = 0; i < N; i++)
+    y[i] = 42;
+  #pragma omp target data map(y) use_device_addr(y)
+    #pragma omp target has_device_addr(y)
+      for (int i = 0; i < N; i++)
+	y[i] = i;
+  for (int i = 0; i < N; i++)
+    if (y[i] != i)
+      __builtin_abort ();
+
+  #pragma omp target data map(y[:N]) use_device_addr(y)
+    #pragma omp target has_device_addr(y[:N])
+      for (int i = 0; i < N; i++)
+	y[i] = i + 2;
+  for (int i = 0; i < N; i++)
+    if (y[i] != i + 2)
+      __builtin_abort ();
+
+  #pragma omp target data map(y[:N]) use_device_addr(y)
+    #pragma omp target has_device_addr(y[24])
+	y[24] = 42;
+  if (y[24] != 42)
+    __builtin_abort ();
+
+  #pragma omp target data map(y[:N]) use_device_addr(y)
+    #pragma omp target has_device_addr(y[24:])
+      for (int i = 24; i < N; i++)
+	y[i] = i + 3;
+  for (int i = 24; i < N; i++)
+    if (y[i] != i + 3)
+      __builtin_abort ();
+
+  #pragma omp target data map(y[:N]) use_device_addr(y)
+    #pragma omp target has_device_addr(y[12:24])
+      for (int i = 12; i < 24; i++)
+	y[i] = i + 4;
+  for (int i = 12; i < 24; i++)
+    if (y[i] != i + 4)
+      __builtin_abort ();
+
+  int u[0];
+  #pragma omp target data map(u) use_device_addr(u)
+    #pragma omp target has_device_addr(u)
+  ;
+
+  struct S { int m; } s;
+  s.m = 42;
+  #pragma omp target data map (s) use_device_addr (s)
+    #pragma omp target has_device_addr (s)
+      ++s.m;
+  if (s.m != 43)
+    __builtin_abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c b/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c
new file mode 100644
index 00000000000..fd99a82f66a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c
@@ -0,0 +1,33 @@
+/* Testing 'has_device_addr' clause with variable sized array. */
+
+int
+foo (int size)
+{
+  int x[size];
+
+  #pragma omp target data map(x[:size]) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      for (int i = 0; i < size; i++)
+	x[i] = i;
+  for (int i = 0; i < size; i++)
+    if (x[i] != i)
+      __builtin_abort ();
+
+  #pragma omp target data map(x) use_device_addr(x)
+    #pragma omp target has_device_addr(x[2:3])
+      for (int i = 0; i < size; i++)
+	x[i] = i;
+  for (int i = 0; i < size; i++)
+    if (x[i] != i)
+      __builtin_abort ();
+
+  return 0;
+}
+
+int
+main ()
+{
+  foo (40);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90
new file mode 100644
index 00000000000..2945864fa53
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90
@@ -0,0 +1,50 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none
+
+  integer, parameter :: N = 40
+  integer :: x, i
+  integer :: y (N)
+  integer :: u (0)
+
+  x = 24
+  !$omp target data map(x) use_device_addr(x)
+    !$omp target has_device_addr(x)
+      x = 42;
+    !$omp end target
+  !$omp end target data
+  if (x /= 42) stop 1
+
+  y = 42
+  !$omp target data map(y) use_device_addr(y)
+    !$omp target has_device_addr(y)
+      y = [(i, i=1, N)]
+    !$omp end target
+  !$omp end target data
+  if (any (y /= [(i, i = 1, N)])) stop 2
+
+  !$omp target data map(y(:N)) use_device_addr(y)
+    !$omp target has_device_addr(y(:N))
+      y = [(i+2, i=1, N)]
+    !$omp end target
+  !$omp end target data
+  if (any (y /= [(i+2, i = 1, N)])) stop 3
+
+  !$omp target data map(y) use_device_addr(y)
+    !$omp target has_device_addr(y(24:))
+      do i = 24, N
+        y(i) = i + 3
+      end do
+    !$omp end target
+  !$omp end target data
+  do i = 24, N
+    if (y(i) /= i + 3) stop 5
+  end do
+
+  !$omp target data map(u) use_device_addr(u)
+    !$omp target has_device_addr(u)
+    !$omp end target
+  !$omp end target data
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90
new file mode 100644
index 00000000000..a8d78a75af3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90
@@ -0,0 +1,40 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none
+
+  integer, parameter :: N = 5
+  integer :: i, x(N), y(N), z(N:2*N-1)
+  target :: z
+
+  x = 42
+  y = 43
+  z = 44
+
+  call foo (x, y, z)
+  if (any (x /= [(i, i = 1, N)])) stop 1
+  if (any (y /= [(2*i, i = 1, N)])) stop 2
+  if (any (z /= [(3*i, i = 1, N)])) stop 3
+
+  contains
+  subroutine foo(a, b, c)
+    integer :: a(:)
+    integer :: b(*)
+    integer, pointer, intent(in) :: c(:)
+
+    !$omp target data map(a,b(:N),c) use_device_addr(a,b(:N),c)
+      !$omp target has_device_addr(A,B(:N),C)
+        if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 10
+        if (lbound(b,dim=1) /= 1) stop 11
+        if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 12
+        if (any (a /= 42)) stop 13
+        if (any (b(:N) /= 43)) stop 14
+        if (any (c /= 44)) stop 15
+        a = [(i, i=1, N)]
+        b(:N) = [(2*i, i = 1, N)]
+        c = [(3*i, i = 1, N)]
+      !$omp end target
+    !$omp end target data
+  end subroutine foo
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90
new file mode 100644
index 00000000000..c6293b4de2e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90
@@ -0,0 +1,90 @@
+! Test optional dummy arguments in HAS_DEVICE_ADDR.
+
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none
+
+  integer, target :: x
+  integer, pointer :: ptr
+  integer, parameter :: N=7
+  real :: y1(N), y2(N)
+  integer, target :: y3(N:2*N-1)
+  integer :: i
+
+  x = 24
+  ptr => x
+  y1 = 42.24
+  y2 = 42.24
+  y3 = 42
+
+  call optional_scalar (is_present=.false.)
+  if (x /= 24) stop 1
+
+  call optional_scalar (x, is_present=.true.)
+  if (x /= 42) stop 2
+
+  call optional_ptr (is_present=.false.)
+  if (x /= 42) stop 3
+  if (ptr /= 42) stop 4
+
+  call optional_ptr (ptr, is_present=.true.)
+  if (x /= 84) stop 5
+  if (ptr /= 84) stop 6
+
+  call optional_array (is_present=.false.)
+  if (any (y1 /= [(42.24, i=1, N)])) stop 7
+  if (any (y2 /= [(42.24, i=1, N)])) stop 8
+  if (any (y3 /= [(42, i=1, N)])) stop 9
+
+  call optional_array (y1, y2, y3, is_present=.true.)
+  if (any (y1 /= [(42.24+i, i=1, N)])) stop 10
+  if (any (y2 /= [(42.24+2*i, i=1, N)])) stop 11
+  if (any (y3 /= [(42+3*i, i=1, N)])) stop 12
+
+contains
+  subroutine optional_scalar (a, is_present)
+    integer, optional :: a
+    logical, value :: is_present
+
+    !$omp target data map(a) use_device_addr(a)
+      !$omp target has_device_addr(a)
+        if (is_present) a = 42
+      !$omp end target
+    !$omp end target data
+  end subroutine optional_scalar
+
+  subroutine optional_ptr (a, is_present)
+    integer, pointer, optional :: a
+    logical, value :: is_present
+    !$omp target data map(a) use_device_addr(a)
+      !$omp target has_device_addr(a)
+        if (is_present) a = 84
+      !$omp end target
+    !$omp end target data
+  end subroutine optional_ptr
+
+  subroutine optional_array (a, b, c, is_present)
+    real, optional :: a(:), b(*)
+    integer, optional, pointer, intent(in) :: c(:)
+    logical, value :: is_present
+    integer :: i
+
+    !$omp target data map(a, b(:N), c) use_device_addr(a, b, c)
+      !$omp target has_device_addr(a, b, c)
+        if (is_present) then
+          if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 21
+          if (lbound(b,dim=1) /= 1) stop 22
+          if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 23
+          if (any (a /= [(42.24, i = 1, N)])) stop 24
+          if (any (b(:N) /= [(42.24, i = 1, N)])) stop 25
+          if (any (c /= [(42, i = 1, N)])) stop 26
+          a = [(42.24+i, i=1, N)]
+          b(:N) = [(42.24+2*i, i=1, N)]
+          c = [(42+3*i, i=1, N)]
+        end if
+      !$omp end target
+    !$omp end target data
+  end subroutine optional_array
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90
new file mode 100644
index 00000000000..59d3e3d31dd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90
@@ -0,0 +1,71 @@
+! Test allocatables in HAS_DEVICE_ADDR.
+
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none
+
+  integer, parameter :: N = 5
+  integer, allocatable :: x
+  integer, allocatable :: y(:)
+  call scalar_dummy (x)
+  call array_dummy (y)
+  call array_dummy_optional (y)
+  call array_dummy_optional ()
+
+contains
+  subroutine scalar_dummy (a)
+    integer, allocatable :: a
+
+    allocate (a)
+    a = 24
+
+    !$omp target data map(a) use_device_addr(a)
+      !$omp target has_device_addr(a)
+        a = 42
+      !$omp end target
+    !$omp end target data
+    if (a /= 42) stop 1
+
+    deallocate (a)
+  end subroutine scalar_dummy
+
+  subroutine array_dummy (a)
+    integer, allocatable :: a(:)
+    integer :: i
+
+    allocate (a(N))
+    a = 42
+
+    !$omp target data map(a) use_device_addr(a)
+      !$omp target has_device_addr(a)
+        a = [(i, i=1, N)]
+      !$omp end target
+    !$omp end target data
+    if (any (a /= [(i, i=1, N)])) stop 2
+
+    deallocate (a)
+  end subroutine array_dummy
+
+  subroutine array_dummy_optional (a)
+    integer, optional, allocatable :: a(:)
+    integer :: i
+
+    if (present (a)) then
+      allocate (a(N))
+      a = 42
+    end if
+
+    !$omp target data map(a) use_device_addr(a)
+      !$omp target has_device_addr(a)
+        if (present (a)) a = [(i, i=1, N)]
+      !$omp end target
+    !$omp end target data
+
+    if (present (a)) then
+      if (any (a /= [(i, i=1, N)])) stop 2
+      deallocate (a)
+    end if
+  end subroutine array_dummy_optional
+
+end program main