OpenMP: Permit in Fortran omp target data without map

gcc/fortran/ChangeLog:

	* openmp.c (resolve_omp_clauses): Permit 'omp target data' without
	map if use_device_{addr,ptr} is present.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/map-3.f90: New test.
	* gfortran.dg/gomp/map-4.f90: New test.

3 files changed, 54 insertions, 3 deletions

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 0fd9988..16f39a4 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5330,17 +5330,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
   if (omp_clauses->depend_source && code->op != EXEC_OMP_ORDERED)
     gfc_error ("SOURCE dependence type only allowed "
 	       "on ORDERED directive at %L", &code->loc);
-  if (!openacc && code && omp_clauses->lists[OMP_LIST_MAP] == NULL)
+  if (!openacc
+      && code
+      && omp_clauses->lists[OMP_LIST_MAP] == NULL
+      && omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR] == NULL
+      && omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR] == NULL)
     {
       const char *p = NULL;
       switch (code->op)
 	{
-	case EXEC_OMP_TARGET_DATA: p = "TARGET DATA"; break;
 	case EXEC_OMP_TARGET_ENTER_DATA: p = "TARGET ENTER DATA"; break;
 	case EXEC_OMP_TARGET_EXIT_DATA: p = "TARGET EXIT DATA"; break;
 	default: break;
 	}
-      if (p)
+      if (code->op == EXEC_OMP_TARGET_DATA)
+	gfc_error ("TARGET DATA must contain at least one MAP, USE_DEVICE_PTR, "
+		   "or USE_DEVICE_ADDR clause at %L", &code->loc);
+      else if (p)
 	gfc_error ("%s must contain at least one MAP clause at %L",
 		   p, &code->loc);
     }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-3.f90
new file mode 100644
index 0000000..13f6364
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-3.f90
@@ -0,0 +1,38 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine bar
+integer, target :: x
+integer, allocatable, target :: y(:,:), z(:,:)
+x = 7
+!$omp target enter data map(to:x)
+
+x = 8
+!$omp target data map(always, to: x)
+call foo(x)
+!$omp end target data
+
+!$omp target data use_device_ptr(x)
+call foo2(x)
+!$omp end target data
+
+!$omp target data use_device_addr(x)
+call foo2(x)
+!$omp end target data
+!$omp target exit data map(release:x)
+
+!$omp target data map(y) use_device_addr(y)
+call foo3(y)
+!$omp end target data
+
+!$omp target data map(z) use_device_ptr(z)
+call foo3(z)
+!$omp end target data
+end 
+
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:x\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,to:x\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_ptr\\(x\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(release:x\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) y.data \\\[len: .*\\) map\\(to:y \\\[pointer set, len: .*\\) map\\(alloc:.*y.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(y\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_ptr\\(z\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-4.f90 b/gcc/testsuite/gfortran.dg/gomp/map-4.f90
new file mode 100644
index 0000000..3aa1c80
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-4.f90
@@ -0,0 +1,7 @@
+!$omp target enter data device(1) if (.true.) nowait ! { dg-error "TARGET ENTER DATA must contain at least one MAP clause" }
+
+!$omp target data device(1)  ! { dg-error "TARGET DATA must contain at least one MAP, USE_DEVICE_PTR, or USE_DEVICE_ADDR clause" }
+!$omp endtarget data
+
+!$omp target exit data device(1) if (.true.) nowait ! { dg-error "TARGET EXIT DATA must contain at least one MAP clause" }
+end