Re: [Patch] OpenMP, libgomp: Add new runtime routine omp_target_is_accessible.

[Marcel Vollweiler <marcel@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 1270 bytes --]

Hi Tobias,

> Minor remark to the test:
>
> On 11.03.22 13:30, Marcel Vollweiler wrote:
>> +  int d = omp_get_default_device ();
> ...
>> +  int shared_mem = 0;
>> +  #pragma omp target map (alloc: shared_mem) device (d)
>> +    shared_mem = 1;
>> +  if (omp_target_is_accessible (p, sizeof (int), d) != shared_mem)
>> +    __builtin_abort ();
>
> I wonder whether it makes sense to do instead
>    for (d = 0; d <= omp_get_num_devices(); ++d)
> instead of just
>    d = omp_get_default_device();
> given that we have already found once in a while bugs when testing more
> than just the default device - be it because devices differed or because
> '0' was special.
>
> In particular, I could image having at the same time two or three devices
> available of type intelmic + gcn + nvptx, possibly mixing shared memory,
> nonshared memory and semi-shared memory*

Good hint, thanks. I updated the C(++) and Fortran tests accordingly and
attached the updated patch.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: target-is-accessible-patch.diff --]
[-- Type: text/plain, Size: 8225 bytes --]

OpenMP, libgomp: Add new runtime routine omp_target_is_accessible.

gcc/ChangeLog:

	* omp-low.cc (omp_runtime_api_call): Added target_is_accessible to
	omp_runtime_apis array.

libgomp/ChangeLog:

	* libgomp.map: Added omp_target_is_accessible.
	* libgomp.texi: Tagged omp_target_is_accessible as supported.
	* omp.h.in: Added omp_target_is_accessible.
	* omp_lib.f90.in: Added interface for omp_target_is_accessible.
	* omp_lib.h.in: Likewise.
	* target.c (omp_target_is_accessible): Added implementation of
	omp_target_is_accessible.
	* testsuite/libgomp.c-c++-common/target-is-accessible-1.c: New test.
	* testsuite/libgomp.fortran/target-is-accessible-1.f90: New test.

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 77176ef..bf38fad 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -3959,6 +3959,7 @@ omp_runtime_api_call (const_tree fndecl)
       "target_associate_ptr",
       "target_disassociate_ptr",
       "target_free",
+      "target_is_accessible",
       "target_is_present",
       "target_memcpy",
       "target_memcpy_rect",
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2ac5809..1764380 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -226,6 +226,11 @@ OMP_5.1 {
 	omp_get_teams_thread_limit_;
 } OMP_5.0.2;
 
+OMP_5.1.1 {
+  global:
+	omp_target_is_accessible;
+} OMP_5.1;
+
 GOMP_1.0 {
   global:
 	GOMP_atomic_end;
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 161a423..58e432c 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -311,7 +311,7 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{omp_set_num_teams}, @code{omp_set_teams_thread_limit},
       @code{omp_get_max_teams}, @code{omp_get_teams_thread_limit} runtime
       routines @tab Y @tab
-@item @code{omp_target_is_accessible} runtime routine @tab N @tab
+@item @code{omp_target_is_accessible} runtime routine @tab Y @tab
 @item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
       runtime routines @tab N @tab
 @item @code{omp_get_mapped_ptr} runtime routine @tab N @tab
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index 89c5d65..1ec7415 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -282,6 +282,8 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
 extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
 				     __SIZE_TYPE__, int) __GOMP_NOTHROW;
 extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
+extern int omp_target_is_accessible (const void *, __SIZE_TYPE__, int)
+  __GOMP_NOTHROW;
 
 extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
 extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index daf40dc..f369507 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -835,6 +835,16 @@
           end function omp_target_disassociate_ptr
         end interface
 
+        interface
+          function omp_target_is_accessible (ptr, size, device_num) bind(c)
+            use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t, c_int
+            integer(c_int) :: omp_target_is_accessible
+            type(c_ptr), value :: ptr
+            integer(c_size_t), value :: size
+            integer(c_int), value :: device_num
+          end function omp_target_is_accessible
+        end interface
+
 #if _OPENMP >= 201811
 !GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
 #endif
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index ff857a4..5ea0366 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -416,3 +416,14 @@
           integer(c_int), value :: device_num
         end function omp_target_disassociate_ptr
       end interface
+
+      interface
+        function omp_target_is_accessible (ptr, size, device_num)          &
+     &      bind(c)
+          use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t, c_int
+          integer(c_int) :: omp_target_is_accessible
+          type(c_ptr), value :: ptr
+          integer(c_size_t), value :: size
+          integer(c_int), value :: device_num
+        end function omp_target_is_accessible
+      end interface
diff --git a/libgomp/target.c b/libgomp/target.c
index 9017458..01d36d1 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3666,6 +3666,24 @@ omp_target_disassociate_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 ())
+    return true;
+
+  struct gomp_device_descr *devicep = resolve_device (device_num);
+  if (devicep == NULL)
+    return false;
+
+  /* TODO: Unified shared memory must be handled when available.  */
+
+  return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM;
+}
+
+int
 omp_pause_resource (omp_pause_resource_t kind, int device_num)
 {
   (void) kind;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
new file mode 100644
index 0000000..7c2cf62
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
@@ -0,0 +1,47 @@
+#include <omp.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int n = omp_get_num_devices ();
+  void *p;
+
+  if (d < 0 || d >= n)
+    d = id;
+
+  if (!omp_target_is_accessible (p, sizeof (int), n))
+    __builtin_abort ();
+
+  if (!omp_target_is_accessible (p, sizeof (int), id))
+    __builtin_abort ();
+
+  if (omp_target_is_accessible (p, sizeof (int), -1))
+    __builtin_abort ();
+
+  if (omp_target_is_accessible (p, sizeof (int), n + 1))
+    __builtin_abort ();
+
+  /* Currently, a host pointer is accessible if the device supports shared
+     memory or omp_target_is_accessible is executed on the host. This
+     test case must be adapted when unified shared memory is avialable.  */
+  int a[128];
+  for (int d = 0; d <= omp_get_num_devices (); d++)
+    {
+      int shared_mem = 0;
+      #pragma omp target map (alloc: shared_mem) device (d)
+	shared_mem = 1;
+      if (omp_target_is_accessible (p, sizeof (int), d) != shared_mem)
+	__builtin_abort ();
+
+      if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem)
+	__builtin_abort ();
+
+      for (int i = 0; i < 128; i++)
+	if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem)
+	  __builtin_abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
new file mode 100644
index 0000000..2611855
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
@@ -0,0 +1,50 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id, n, shared_mem, i
+  integer, target :: a(1:128)
+  type(c_ptr) :: p
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+  n = omp_get_num_devices ()
+
+  if (d < 0 .or. d >= n) &
+    d = id
+
+  if (omp_target_is_accessible (p, c_sizeof (d), n) /= 1) &
+    stop 1
+
+  if (omp_target_is_accessible (p, c_sizeof (d), id) /= 1) &
+    stop 2
+
+  if (omp_target_is_accessible (p, c_sizeof (d), -1) /= 0) &
+    stop 3
+
+  if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) &
+    stop 4
+
+  ! Currently, a host pointer is accessible if the device supports shared
+  ! memory or omp_target_is_accessible is executed on the host. This
+  ! test case must be adapted when unified shared memory is avialable.
+  do d = 0, omp_get_num_devices ()
+    shared_mem = 0;
+    !$omp target map (alloc: shared_mem) device (d)
+      shared_mem = 1;
+    !$omp end target
+
+    if (omp_target_is_accessible (p, c_sizeof (d), d) /= shared_mem) &
+      stop 5;
+
+    if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) &
+      stop 6;
+
+    do i = 1, 128
+      if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) &
+        stop 7;
+    end do
+
+  end do
+
+end program main