OpenMP: Cleanups related to the 'present' modifier

Reduce number of enum values passed to libgomp as
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} have the same semantic as
GOMP_MAP_FORCE_PRESENT (i.e. abort if not present, otherwise ignore);
that's different to GOMP_MAP_ALWAYS_PRESENT_{TO,TOFROM,FROM} which also
abort if not present but copy data when present. This is is a follow-up to
the commit r14-1579-g4ede915d5dde93 done 6 days ago.

Additionally, the commit improves a libgomp run-time and a C/C++ compile-time
error wording and extends testcases a tiny bit.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_clause_map): Reword error message for
	clearness especially with 'omp target (enter/exit) data.'

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_omp_clause_map): Reword error message for
	clearness especially with 'omp target (enter/exit) data.'
	* semantics.cc (handle_omp_array_sections): Handle
	GOMP_MAP_{ALWAYS_,}PRESENT_{TO,TOFROM,FROM,ALLOC} enum values.

gcc/ChangeLog:

	* gimplify.cc (gimplify_adjust_omp_clauses_1): Use
	GOMP_MAP_FORCE_PRESENT for 'present alloc' implicit mapping.
	(gimplify_adjust_omp_clauses): Change
	GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to the equivalent
	GOMP_MAP_FORCE_PRESENT.
	* omp-low.cc (lower_omp_target): Remove handling of no-longer valid
	GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC}; update map kinds used for
	to/from clauses with present modifier.

include/ChangeLog:

	* gomp-constants.h (enum gomp_map_kind): Change the enum values
	GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to be compiler only.
	(GOMP_MAP_PRESENT_P): Update to include also GOMP_MAP_FORCE_PRESENT.

libgomp/ChangeLog:

	* target.c (gomp_to_device_kind_p, gomp_map_vars_internal): Replace
	GOMP_MAP_PRESENT_{FROM,TO,TOFROM,ACLLOC} by GOMP_MAP_FORCE_PRESENT.
	(gomp_map_vars_internal, gomp_update): Likewise; unify and improve
	error message.
	* testsuite/libgomp.c-c++-common/target-present-2.c: Update for
	changed error message.
	* testsuite/libgomp.fortran/target-present-1.f90: Likewise.
	* testsuite/libgomp.fortran/target-present-2.f90: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
	* testsuite/libgomp.c-c++-common/target-present-1.c: Likewise and
	extend testcase to check that data is copied when needed.
	* testsuite/libgomp.c-c++-common/target-present-3.c: Likewise.
	* testsuite/libgomp.fortran/target-present-3.f90: Likewise.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/defaultmap-4.c: Update scan-tree-dump.
	* c-c++-common/gomp/map-9.c: Likewise.
	* gfortran.dg/gomp/defaultmap-8.f90: Likewise.
	* gfortran.dg/gomp/map-11.f90: Likewise.
	* gfortran.dg/gomp/target-update-1.f90: Likewise.
	* gfortran.dg/gomp/map-12.f90: Likewise; also check original dump.
	* c-c++-common/gomp/map-6.c: Update dg-error and also check
	clause error with 'target (enter/exit) data'.

8 files changed, 75 insertions, 42 deletions

diff --git a/libgomp/target.c b/libgomp/target.c
index a9e8005..e3c4121 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -358,8 +358,8 @@ gomp_to_device_kind_p (int kind)
     case GOMP_MAP_FORCE_ALLOC:
     case GOMP_MAP_FORCE_FROM:
     case GOMP_MAP_ALWAYS_FROM:
-    case GOMP_MAP_PRESENT_FROM:
     case GOMP_MAP_ALWAYS_PRESENT_FROM:
+    case GOMP_MAP_FORCE_PRESENT:
       return false;
     default:
       return true;
@@ -1699,37 +1699,29 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		    i = j - 1;
 		    break;
 		  case GOMP_MAP_FORCE_PRESENT:
+		  case GOMP_MAP_ALWAYS_PRESENT_TO:
+		  case GOMP_MAP_ALWAYS_PRESENT_FROM:
+		  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 		    {
 		      /* We already looked up the memory region above and it
 			 was missing.  */
 		      size_t size = k->host_end - k->host_start;
 		      gomp_mutex_unlock (&devicep->lock);
 #ifdef HAVE_INTTYPES_H
-		      gomp_fatal ("present clause: !acc_is_present (%p, "
-				  "%"PRIu64" (0x%"PRIx64"))",
-				  (void *) k->host_start,
-				  (uint64_t) size, (uint64_t) size);
+		      gomp_fatal ("present clause: not present on the device "
+				  "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), "
+				  "dev: %d)", (void *) k->host_start,
+				  (uint64_t) size, (uint64_t) size,
+				  devicep->target_id);
 #else
-		      gomp_fatal ("present clause: !acc_is_present (%p, "
-				  "%lu (0x%lx))", (void *) k->host_start,
-				  (unsigned long) size, (unsigned long) size);
+		      gomp_fatal ("present clause: not present on the device "
+				  "(addr: %p, size: %lu (0x%lx), dev: %d)",
+				  (void *) k->host_start,
+				  (unsigned long) size, (unsigned long) size,
+				  devicep->target_id);
 #endif
 		    }
 		    break;
-		  case GOMP_MAP_PRESENT_ALLOC:
-		  case GOMP_MAP_PRESENT_TO:
-		  case GOMP_MAP_PRESENT_FROM:
-		  case GOMP_MAP_PRESENT_TOFROM:
-		  case GOMP_MAP_ALWAYS_PRESENT_TO:
-		  case GOMP_MAP_ALWAYS_PRESENT_FROM:
-		  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
-		    /* We already looked up the memory region above and it
-		       was missing.  */
-		    gomp_mutex_unlock (&devicep->lock);
-		    gomp_fatal ("present clause: not present on the device "
-				"(%p, %d)",
-				(void *) k->host_start, devicep->target_id);
-		    break;
 		  case GOMP_MAP_FORCE_DEVICEPTR:
 		    assert (k->host_end - k->host_start == sizeof (void *));
 		    gomp_copy_host2dev (devicep, aq,
@@ -2149,9 +2141,18 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 		/* We already looked up the memory region above and it
 		   was missing.  */
 		gomp_mutex_unlock (&devicep->lock);
+#ifdef HAVE_INTTYPES_H
 		gomp_fatal ("present clause: not present on the device "
-			    "(%p, %d)",
-			    (void *) hostaddrs[i], devicep->target_id);
+			    "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), "
+			    "dev: %d)", (void *) hostaddrs[i],
+			    (uint64_t) sizes[i], (uint64_t) sizes[i],
+			    devicep->target_id);
+#else
+		gomp_fatal ("present clause: not present on the device "
+			    "(addr: %p, size: %lu (0x%lx), dev: %d)",
+			    (void *) hostaddrs[i], (unsigned long) sizes[i],
+			    (unsigned long) sizes[i], devicep->target_id);
+#endif
 	      }
 	  }
       }
@@ -3465,9 +3466,7 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_FORCE_TOFROM:
 	      case GOMP_MAP_ALWAYS_TO:
 	      case GOMP_MAP_ALWAYS_TOFROM:
-	      case GOMP_MAP_PRESENT_FROM:
-	      case GOMP_MAP_PRESENT_TO:
-	      case GOMP_MAP_PRESENT_TOFROM:
+	      case GOMP_MAP_FORCE_PRESENT:
 	      case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	      case GOMP_MAP_ALWAYS_PRESENT_TO:
 	      case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
@@ -3710,8 +3709,6 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_FORCE_TOFROM:
 	      case GOMP_MAP_ALWAYS_FROM:
 	      case GOMP_MAP_ALWAYS_TOFROM:
-	      case GOMP_MAP_PRESENT_FROM:
-	      case GOMP_MAP_PRESENT_TOFROM:
 	      case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	      case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 		copy = true;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
index aa34319..5eaa9cd 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
@@ -4,24 +4,34 @@
 
 int main (void)
 {
-  int a[N], b[N], c[N];
+  int a[N], b[N], c[N], d[N];
 
   for (int i = 0; i < N; i++) {
     a[i] = i * 2;
     b[i] = i * 3 + 1;
+    d[i] = i * 5;
   }
 
-  #pragma omp target enter data map (alloc: a, c)
-    /* a has already been allocated, so this should be okay.  */
-    #pragma omp target map (present, to: a)
+  #pragma omp target enter data map (alloc: c, d) map(to: a)
+    #pragma omp target map (present, always, to: d)
+      for (int i = 0; i < N; i++)
+	if (d[i] != i * 5)
+	  __builtin_abort ();
+
+    /* a has already been mapped and 'c' allocated so this should be okay.  */
+    #pragma omp target map (present, to: a) map(present, always, from: c)
       for (int i = 0; i < N; i++)
 	c[i] = a[i];
 
+    for (int i = 0; i < N; i++)
+      if (c[i] != i * 2)
+	__builtin_abort ();
+
     fprintf (stderr, "CheCKpOInT\n");
     /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
 
     /* b has not been allocated, so this should result in an error.  */
-    /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
+    /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
     /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
     #pragma omp target map (present, to: b)
       for (int i = 0; i < N; i++)
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
index ad11023..07ae90b 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
@@ -21,7 +21,7 @@ int main (void)
     /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
 
     /* b has not been allocated, so this should result in an error.  */
-    /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
+    /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
     /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
     #pragma omp target defaultmap (present)
       for (int i = 0; i < N; i++)
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
index 455519a..582247d 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
@@ -16,11 +16,24 @@ int main (void)
   /* This should work as a has already been allocated.  */
   #pragma omp target update to (present: a)
 
+  #pragma omp target map(present,alloc: a, c)
+  for (int i = 0; i < N; i++) {
+    if (a[i] != i * 2)
+      __builtin_abort ();
+    c[i] = 23*i;
+  }
+
+  #pragma omp target update from(present : c)
+  for (int i = 0; i < N; i++) {
+    if (c[i] != 23*i)
+      __builtin_abort ();
+  }
+
   fprintf (stderr, "CheCKpOInT\n");
   /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
 
   /* This should fail as b has not been allocated.  */
-  /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
+    /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
   /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
   #pragma omp target update to (present: b)
 
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
index 768166f..fc13609 100644
--- a/libgomp/testsuite/libgomp.fortran/target-present-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
@@ -20,7 +20,7 @@ program main
     ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
 
     ! b has not been allocated, so this should result in an error.
-    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } }
+    ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } }
     ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } }
     !$omp target map (present, to: b)
       do i = 1, N
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
index 8f2c24e..524d01d 100644
--- a/libgomp/testsuite/libgomp.fortran/target-present-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
@@ -20,7 +20,7 @@ program main
     ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
 
     ! b has not been allocated, so this should result in an error.
-    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } }
+    ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } }
     ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } }
     !$omp target defaultmap (present)
       do i = 1, N
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
index eb29c90..dd4af4c 100644
--- a/libgomp/testsuite/libgomp.fortran/target-present-3.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
@@ -9,14 +9,27 @@ program main
   end do
 
   !$omp target enter data map (alloc: a, c)
-    ! This should work as a has already been allocated.
-    !$omp target update to (present: a)
+
+  ! This should work as a has already been allocated.
+  !$omp target update to (present: a)
+
+  !$omp target map(present, alloc: a, c)
+    do i = 1, N
+      if (a(i) /= i * 2) stop 1
+      c(i) = 23 * i
+    end do
+  !$omp end target
+
+  !$omp target update from (present: c)
+    do i = 1, N
+      if (c(i) /= 23 * i) stop 1
+    end do
 
     print *, "CheCKpOInT"
     ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
 
     ! This should fail as b has not been allocated.
-    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } }
+    ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } }
     ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } }
     !$omp target update to (present: b)
   !$omp target exit data map (from: c)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c
index 61c8109..02fbfda 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c
@@ -48,5 +48,5 @@ main (int argc, char **argv)
 }
 
 /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "present clause: !acc_is_present" } */
+/* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" } */
 /* { dg-shouldfail "" } */