openmp: Conforming device numbers and omp_{initial,invalid}_device

OpenMP 5.2 changed once more what device numbers are allowed.
In 5.1, valid device numbers were [0, omp_get_num_devices()].
5.2 makes also -1 valid (calls it omp_initial_device), which is equivalent
in behavior to omp_get_num_devices() number but has the advantage that it
is a constant.  And it also introduces omp_invalid_device which is
also a constant with implementation defined value < -1.  That value should
act like sNaN, any time any device construct (GOMP_target*) or OpenMP runtime
API routine is asked for such a device, the program is terminated.
And if OMP_TARGET_OFFLOAD=mandatory, all non-conforming device numbers (which
is all but [-1, omp_get_num_devices()] other than omp_invalid_device)
must be treated like omp_invalid_device.

For device constructs, we have a compatibility problem, we've historically
used 2 magic negative values to mean something special.
GOMP_DEVICE_ICV (-1) means device clause wasn't present, pick the
		     omp_get_default_device () number
GOMP_DEVICE_FALLBACK (-2) means the host device (this is used e.g. for
			  #pragma omp target if (cond)
			  where if cond is false, we pass -2
But 5.2 requires that omp_initial_device is -1 (there were discussions
about it, advantage of -1 is that one can say iterate over the
[-1, omp_get_num_devices()-1] range to get all devices starting with
the host/initial one.
And also, if user passes -2, unless it is omp_invalid_device, we need to
treat it like non-conforming with OMP_TARGET_OFFLOAD=mandatory.

So, the patch does on the compiler side some number remapping,
user_device_num >= -2U ? user_device_num - 1 : user_device_num.
This remapping is done at compile time if device clause has constant
argument, otherwise at runtime, and means that for user -1 (omp_initial_device)
we pass -2 to GOMP_* in the runtime library where it treats it like host
fallback, while -2 is remapped to -3 (one of the non-conforming device numbers,
for those it doesn't matter which one is which).
omp_invalid_device is then -4.
For the OpenMP device runtime APIs, no remapping is done.

This patch doesn't deal with the initial default-device-var for
OMP_TARGET_OFFLOAD=mandatory , the spec says that the inital ICV value
for that should in that case depend on whether there are any offloading
devices or not (if not, should be omp_invalid_device), but that means
we can't determine the number of devices lazily (and let libraries have the
possibility to register their offloading data etc.).

2022-06-13  Jakub Jelinek  <jakub@redhat.com>

gcc/
	* omp-expand.cc (expand_omp_target): Remap user provided
	device clause arguments, -1 to -2 and -2 to -3, either
	at compile time if constant, or at runtime.
include/
	* gomp-constants.h (GOMP_DEVICE_INVALID): Define.
libgomp/
	* omp.h.in (omp_initial_device, omp_invalid_device): New enumerators.
	* omp_lib.f90.in (omp_initial_device, omp_invalid_device): New
	parameters.
	* omp_lib.h.in (omp_initial_device, omp_invalid_device): Likewise.
	* target.c (resolve_device): Add remapped argument, handle
	GOMP_DEVICE_ICV only if remapped is true (and clear remapped),
	for negative values, treat GOMP_DEVICE_FALLBACK as fallback only
	if remapped, otherwise treat omp_initial_device that way.  For
	omp_invalid_device, always emit gomp_fatal, even when
	OMP_TARGET_OFFLOAD isn't mandatory.
	(GOMP_target, GOMP_target_ext, GOMP_target_data, GOMP_target_data_ext,
	GOMP_target_update, GOMP_target_update_ext,
	GOMP_target_enter_exit_data): Pass true as remapped argument to
	resolve_device.
	(omp_target_alloc, omp_target_free, omp_target_is_present,
	omp_target_memcpy_check, omp_target_associate_ptr,
	omp_target_disassociate_ptr, omp_get_mapped_ptr,
	omp_target_is_accessible): Pass false as remapped argument to
	resolve_device.  Treat omp_initial_device the same as
	gomp_get_num_devices ().  Don't bypass resolve_device calls if
	device_num is negative.
	(omp_pause_resource): Treat omp_initial_device the same as
	gomp_get_num_devices ().  Call resolve_device.
	* icv-device.c (omp_set_default_device): Always set to device_num
	even when it is negative.
	* libgomp.texi: Document that Conforming device numbers,
	omp_initial_device and omp_invalid_device is implemented.
	* testsuite/libgomp.c/target-41.c (main): Add test with
	omp_initial_device.
	* testsuite/libgomp.c/target-45.c: New test.
	* testsuite/libgomp.c/target-46.c: New test.
	* testsuite/libgomp.c/target-47.c: New test.
	* testsuite/libgomp.c-c++-common/target-is-accessible-1.c (main): Add
	test with omp_initial_device.  Use -5 instead of -1 for negative value
	test.
	* testsuite/libgomp.fortran/target-is-accessible-1.f90 (main):
	Likewise.  Reorder stop numbers.

1 files changed, 60 insertions, 67 deletions

diff --git a/libgomp/target.c b/libgomp/target.c
index 4740f8a..c0844f2 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -126,18 +126,31 @@ gomp_get_num_devices (void)
 }
 
 static struct gomp_device_descr *
-resolve_device (int device_id)
+resolve_device (int device_id, bool remapped)
 {
-  if (device_id == GOMP_DEVICE_ICV)
+  if (remapped && device_id == GOMP_DEVICE_ICV)
     {
       struct gomp_task_icv *icv = gomp_icv (false);
       device_id = icv->default_device_var;
+      remapped = false;
     }
 
-  if (device_id < 0 || device_id >= gomp_get_num_devices ())
+  if (device_id < 0)
+    {
+      if (device_id == (remapped ? GOMP_DEVICE_HOST_FALLBACK
+				 : omp_initial_device))
+	return NULL;
+      if (device_id == omp_invalid_device)
+	gomp_fatal ("omp_invalid_device encountered");
+      else if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY)
+	gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
+		    "but device not found");
+
+      return NULL;
+    }
+  else if (device_id >= gomp_get_num_devices ())
     {
       if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
-	  && device_id != GOMP_DEVICE_HOST_FALLBACK
 	  && device_id != num_devices_openmp)
 	gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
 		    "but device not found");
@@ -2588,7 +2601,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
 	     size_t mapnum, void **hostaddrs, size_t *sizes,
 	     unsigned char *kinds)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
 
   void *fn_addr;
   if (devicep == NULL
@@ -2647,7 +2660,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
 		 void **hostaddrs, size_t *sizes, unsigned short *kinds,
 		 unsigned int flags, void **depend, void **args)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
   size_t tgt_align = 0, tgt_size = 0;
   bool fpc_done = false;
 
@@ -2805,7 +2818,7 @@ void
 GOMP_target_data (int device, const void *unused, size_t mapnum,
 		  void **hostaddrs, size_t *sizes, unsigned char *kinds)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
 
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
@@ -2824,7 +2837,7 @@ void
 GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
 		      size_t *sizes, unsigned short *kinds)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
 
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
@@ -2855,7 +2868,7 @@ void
 GOMP_target_update (int device, const void *unused, size_t mapnum,
 		    void **hostaddrs, size_t *sizes, unsigned char *kinds)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
 
   if (devicep == NULL
       || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
@@ -2870,7 +2883,7 @@ GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
 			size_t *sizes, unsigned short *kinds,
 			unsigned int flags, void **depend)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
 
   /* If there are depend clauses, but nowait is not present,
      block the parent task until the dependencies are resolved
@@ -3063,7 +3076,7 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
 			     size_t *sizes, unsigned short *kinds,
 			     unsigned int flags, void **depend)
 {
-  struct gomp_device_descr *devicep = resolve_device (device);
+  struct gomp_device_descr *devicep = resolve_device (device, true);
 
   /* If there are depend clauses, but nowait is not present,
      block the parent task until the dependencies are resolved
@@ -3296,13 +3309,11 @@ GOMP_teams4 (unsigned int num_teams_low, unsigned int num_teams_high,
 void *
 omp_target_alloc (size_t size, int device_num)
 {
-  if (device_num == gomp_get_num_devices ())
+  if (device_num == omp_initial_device
+      || device_num == gomp_get_num_devices ())
     return malloc (size);
 
-  if (device_num < 0)
-    return NULL;
-
-  struct gomp_device_descr *devicep = resolve_device (device_num);
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
   if (devicep == NULL)
     return NULL;
 
@@ -3319,20 +3330,15 @@ omp_target_alloc (size_t size, int device_num)
 void
 omp_target_free (void *device_ptr, int device_num)
 {
-  if (device_ptr == NULL)
-    return;
-
-  if (device_num == gomp_get_num_devices ())
+  if (device_num == omp_initial_device
+      || device_num == gomp_get_num_devices ())
     {
       free (device_ptr);
       return;
     }
 
-  if (device_num < 0)
-    return;
-
-  struct gomp_device_descr *devicep = resolve_device (device_num);
-  if (devicep == NULL)
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
+  if (devicep == NULL || device_ptr == NULL)
     return;
 
   if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
@@ -3350,19 +3356,17 @@ omp_target_free (void *device_ptr, int device_num)
 int
 omp_target_is_present (const void *ptr, int device_num)
 {
-  if (ptr == NULL)
+  if (device_num == omp_initial_device
+      || device_num == gomp_get_num_devices ())
     return 1;
 
-  if (device_num == gomp_get_num_devices ())
-    return 1;
-
-  if (device_num < 0)
-    return 0;
-
-  struct gomp_device_descr *devicep = resolve_device (device_num);
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
   if (devicep == NULL)
     return 0;
 
+  if (ptr == NULL)
+    return 1;
+
   if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
       || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
     return 1;
@@ -3384,12 +3388,11 @@ omp_target_memcpy_check (int dst_device_num, int src_device_num,
 			 struct gomp_device_descr **dst_devicep,
 			 struct gomp_device_descr **src_devicep)
 {
-  if (dst_device_num != gomp_get_num_devices ())
+  if (dst_device_num != gomp_get_num_devices ()
+      /* Above gomp_get_num_devices has to be called unconditionally.  */
+      && dst_device_num != omp_initial_device)
     {
-      if (dst_device_num < 0)
-	return EINVAL;
-
-      *dst_devicep = resolve_device (dst_device_num);
+      *dst_devicep = resolve_device (dst_device_num, false);
       if (*dst_devicep == NULL)
 	return EINVAL;
 
@@ -3398,12 +3401,10 @@ omp_target_memcpy_check (int dst_device_num, int src_device_num,
 	*dst_devicep = NULL;
     }
 
-  if (src_device_num != num_devices_openmp)
+  if (src_device_num != num_devices_openmp
+      && src_device_num != omp_initial_device)
     {
-      if (src_device_num < 0)
-	return EINVAL;
-
-      *src_devicep = resolve_device (src_device_num);
+      *src_devicep = resolve_device (src_device_num, false);
       if (*src_devicep == NULL)
 	return EINVAL;
 
@@ -3767,13 +3768,11 @@ int
 omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
 			  size_t size, size_t device_offset, int device_num)
 {
-  if (device_num == gomp_get_num_devices ())
-    return EINVAL;
-
-  if (device_num < 0)
+  if (device_num == omp_initial_device
+      || device_num == gomp_get_num_devices ())
     return EINVAL;
 
-  struct gomp_device_descr *devicep = resolve_device (device_num);
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
   if (devicep == NULL)
     return EINVAL;
 
@@ -3830,13 +3829,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
 int
 omp_target_disassociate_ptr (const void *ptr, int device_num)
 {
-  if (device_num == gomp_get_num_devices ())
-    return EINVAL;
-
-  if (device_num < 0)
-    return EINVAL;
-
-  struct gomp_device_descr *devicep = resolve_device (device_num);
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
   if (devicep == NULL)
     return EINVAL;
 
@@ -3872,13 +3865,11 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
 void *
 omp_get_mapped_ptr (const void *ptr, int device_num)
 {
-  if (device_num < 0 || device_num > gomp_get_num_devices ())
-    return NULL;
-
-  if (device_num == omp_get_initial_device ())
+  if (device_num == omp_initial_device
+      || device_num == omp_get_initial_device ())
     return (void *) ptr;
 
-  struct gomp_device_descr *devicep = resolve_device (device_num);
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
   if (devicep == NULL)
     return NULL;
 
@@ -3910,13 +3901,11 @@ omp_get_mapped_ptr (const void *ptr, int device_num)
 int
 omp_target_is_accessible (const void *ptr, size_t size, int device_num)
 {
-  if (device_num < 0 || device_num > gomp_get_num_devices ())
-    return false;
-
-  if (device_num == gomp_get_num_devices ())
+  if (device_num == omp_initial_device
+      || device_num == gomp_get_num_devices ())
     return true;
 
-  struct gomp_device_descr *devicep = resolve_device (device_num);
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
   if (devicep == NULL)
     return false;
 
@@ -3929,10 +3918,14 @@ int
 omp_pause_resource (omp_pause_resource_t kind, int device_num)
 {
   (void) kind;
-  if (device_num == gomp_get_num_devices ())
+  if (device_num == omp_initial_device
+      || device_num == gomp_get_num_devices ())
     return gomp_pause_host ();
-  if (device_num < 0 || device_num >= num_devices_openmp)
+
+  struct gomp_device_descr *devicep = resolve_device (device_num, false);
+  if (devicep == NULL)
     return -1;
+
   /* Do nothing for target devices for now.  */
   return 0;
 }