diff mbox series

OpenMP: Permit in Fortran omp target data without map

Message ID 24e69703-ad51-3c06-66d1-cacea7728abc@codesourcery.com
State New
Headers show
Series OpenMP: Permit in Fortran omp target data without map | expand

Commit Message

Tobias Burnus July 29, 2020, 12:57 p.m. UTC 
I have missed this OpenMP 5 feature already a couple of
times myself ...
(The other related OpenMP 5 feature is that 'map(x) use_device_ptr(x)'
is permitted on the same directive; but that already works and
is tested for in the large use_device*.f90 libgomp tests.)

OK?

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Comments

Jakub Jelinek July 29, 2020, 1:02 p.m. UTC | #1 
On Wed, Jul 29, 2020 at 02:57:25PM +0200, Tobias Burnus wrote:
> I have missed this OpenMP 5 feature already a couple of
> times myself ...
> (The other related OpenMP 5 feature is that 'map(x) use_device_ptr(x)'
> is permitted on the same directive; but that already works and
> is tested for in the large use_device*.f90 libgomp tests.)
> 
> OK?

Ok, thanks.

	Jakub
diff mbox series

Patch

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.

 gcc/fortran/openmp.c                     | 12 +++++++---
 gcc/testsuite/gfortran.dg/gomp/map-3.f90 | 38 ++++++++++++++++++++++++++++++++
 gcc/testsuite/gfortran.dg/gomp/map-4.f90 |  7 ++++++
 3 files changed, 54 insertions(+), 3 deletions(-)

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 0fd998839b2..16f39a4e086 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 00000000000..13f63647bda
--- /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 00000000000..3aa1c8096d7
--- /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