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, 64 insertions, 10 deletions

diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 0821b8d..e7a8af4 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -9983,6 +9983,8 @@ expand_omp_target (struct omp_region *region)
   tree device = NULL_TREE;
   location_t device_loc = UNKNOWN_LOCATION;
   tree goacc_flags = NULL_TREE;
+  bool need_device_adjustment = false;
+  gimple_stmt_iterator adj_gsi;
   if (is_gimple_omp_oacc (entry_stmt))
     {
       /* By default, no GOACC_FLAGs are set.  */
@@ -9994,6 +9996,19 @@ expand_omp_target (struct omp_region *region)
       if (c)
 	{
 	  device = OMP_CLAUSE_DEVICE_ID (c);
+	  /* Ensure 'device' is of the correct type.  */
+	  device = fold_convert_loc (device_loc, integer_type_node, device);
+	  if (TREE_CODE (device) == INTEGER_CST)
+	    {
+	      if (wi::to_wide (device) == GOMP_DEVICE_ICV)
+		device = build_int_cst (integer_type_node,
+					GOMP_DEVICE_HOST_FALLBACK);
+	      else if (wi::to_wide (device) == GOMP_DEVICE_HOST_FALLBACK)
+		device = build_int_cst (integer_type_node,
+					GOMP_DEVICE_HOST_FALLBACK - 1);
+	    }
+	  else
+	    need_device_adjustment = true;
 	  device_loc = OMP_CLAUSE_LOCATION (c);
 	  if (OMP_CLAUSE_DEVICE_ANCESTOR (c))
 	    sorry_at (device_loc, "%<ancestor%> not yet supported");
@@ -10021,7 +10036,8 @@ expand_omp_target (struct omp_region *region)
   if (c)
     cond = OMP_CLAUSE_IF_EXPR (c);
   /* If we found the clause 'if (cond)', build:
-     OpenACC: goacc_flags = (cond ? goacc_flags : flags | GOACC_FLAG_HOST_FALLBACK)
+     OpenACC: goacc_flags = (cond ? goacc_flags
+				  : goacc_flags | GOACC_FLAG_HOST_FALLBACK)
      OpenMP: device = (cond ? device : GOMP_DEVICE_HOST_FALLBACK) */
   if (cond)
     {
@@ -10029,20 +10045,13 @@ expand_omp_target (struct omp_region *region)
       if (is_gimple_omp_oacc (entry_stmt))
 	tp = &goacc_flags;
       else
-	{
-	  /* Ensure 'device' is of the correct type.  */
-	  device = fold_convert_loc (device_loc, integer_type_node, device);
-
-	  tp = &device;
-	}
+	tp = &device;
 
       cond = gimple_boolify (cond);
 
       basic_block cond_bb, then_bb, else_bb;
       edge e;
-      tree tmp_var;
-
-      tmp_var = create_tmp_var (TREE_TYPE (*tp));
+      tree tmp_var = create_tmp_var (TREE_TYPE (*tp));
       if (offloaded)
 	e = split_block_after_labels (new_bb);
       else
@@ -10067,6 +10076,7 @@ expand_omp_target (struct omp_region *region)
       gsi = gsi_start_bb (then_bb);
       stmt = gimple_build_assign (tmp_var, *tp);
       gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+      adj_gsi = gsi;
 
       gsi = gsi_start_bb (else_bb);
       if (is_gimple_omp_oacc (entry_stmt))
@@ -10099,6 +10109,50 @@ expand_omp_target (struct omp_region *region)
       if (device != NULL_TREE)
 	device = force_gimple_operand_gsi (&gsi, device, true, NULL_TREE,
 					   true, GSI_SAME_STMT);
+      if (need_device_adjustment)
+	{
+	  tree tmp_var = create_tmp_var (TREE_TYPE (device));
+	  stmt = gimple_build_assign (tmp_var, device);
+	  gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+	  adj_gsi = gsi_for_stmt (stmt);
+	  device = tmp_var;
+	}
+    }
+
+  if (need_device_adjustment)
+    {
+      tree uns = fold_convert (unsigned_type_node, device);
+      uns = force_gimple_operand_gsi (&adj_gsi, uns, true, NULL_TREE,
+				      false, GSI_CONTINUE_LINKING);
+      edge e = split_block (gsi_bb (adj_gsi), gsi_stmt (adj_gsi));
+      basic_block cond_bb = e->src;
+      basic_block else_bb = e->dest;
+      if (gsi_bb (adj_gsi) == new_bb)
+	{
+	  new_bb = else_bb;
+	  gsi = gsi_last_nondebug_bb (new_bb);
+	}
+
+      basic_block then_bb = create_empty_bb (cond_bb);
+      set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+
+      cond = build2 (GT_EXPR, boolean_type_node, uns,
+		     build_int_cst (unsigned_type_node,
+				    GOMP_DEVICE_HOST_FALLBACK - 1));
+      stmt = gimple_build_cond_empty (cond);
+      adj_gsi = gsi_last_bb (cond_bb);
+      gsi_insert_after (&adj_gsi, stmt, GSI_CONTINUE_LINKING);
+
+      adj_gsi = gsi_start_bb (then_bb);
+      tree add = build2 (PLUS_EXPR, integer_type_node, device,
+			 build_int_cst (integer_type_node, -1));
+      stmt = gimple_build_assign (device, add);
+      gsi_insert_after (&adj_gsi, stmt, GSI_CONTINUE_LINKING);
+
+      make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+      e->flags = EDGE_FALSE_VALUE;
+      add_bb_to_loop (then_bb, cond_bb->loop_father);
+      make_edge (then_bb, else_bb, EDGE_FALLTHRU);
     }
 
   t = gimple_omp_target_data_arg (entry_stmt);