@@ -3909,9 +3909,13 @@ omp_target_is_accessible (const void *ptr, size_t size, int device_num)
if (devicep == NULL)
return false;
- /* TODO: Unified shared memory must be handled when available. */
+ if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return true;
- return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM;
+ if (devicep->is_usm_ptr_func && devicep->is_usm_ptr_func ((void *) ptr))
+ return true;
+
+ return false;
}
int
@@ -23,23 +23,28 @@ main ()
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++)
{
+ /* SHARED_MEM is 1 if and only if host and device share the same memory.
+ OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory. */
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)
+
+ if (shared_mem && !omp_target_is_accessible (p, sizeof (int), d))
+ __builtin_abort ();
+
+ /* USM is disabled by default. Hence OMP_TARGET_IS_ACCESSIBLE should
+ return 0 if shared_mem is false. */
+ if (!shared_mem && omp_target_is_accessible (p, sizeof (int), d))
__builtin_abort ();
- if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem)
+ if (shared_mem && !omp_target_is_accessible (a, 128 * sizeof (int), d))
__builtin_abort ();
for (int i = 0; i < 128; i++)
- if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem)
+ if (shared_mem && !omp_target_is_accessible (&a[i], sizeof (int), d))
__builtin_abort ();
}
new file mode 100644
@@ -0,0 +1,22 @@
+/* { dg-do run } */
+/* { dg-skip-if "USM is only implemented for nvptx." { ! offload_target_nvptx } } */
+
+#include <omp.h>
+#include <stdint.h>
+
+#pragma omp requires unified_shared_memory
+
+int
+main ()
+{
+ int *a = (int *) omp_alloc (sizeof(int), ompx_unified_shared_mem_alloc);
+ if (!a)
+ __builtin_abort ();
+
+ for (int d = 0; d <= omp_get_num_devices (); d++)
+ if (!omp_target_is_accessible (a, sizeof (int), d))
+ __builtin_abort ();
+
+ omp_free(a, ompx_unified_shared_mem_alloc);
+ return 0;
+}
@@ -1,3 +1,5 @@
+! { dg-do run }
+
program main
use omp_lib
use iso_c_binding
@@ -25,24 +27,28 @@ program main
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 is 1 if and only if host and device share the same memory.
+ ! OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory.
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) &
+ if (shared_mem == 1 .and. omp_target_is_accessible (p, c_sizeof (d), d) == 0) &
stop 5;
- if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) &
+ ! USM is disabled by default. Hence OMP_TARGET_IS_ACCESSIBLE should
+ ! return 0 if shared_mem is false.
+ if (shared_mem == 0 .and. omp_target_is_accessible (p, c_sizeof (d), d) /= 0) &
stop 6;
+ if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) == 0) &
+ stop 7;
+
do i = 1, 128
- if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) &
- stop 7;
+ if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) == 0) &
+ stop 8;
end do
end do
new file mode 100644
@@ -0,0 +1,20 @@
+! { dg-do run }
+! { dg-skip-if "USM is only implemented for nvptx." { ! offload_target_nvptx } }
+
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none (external, type)
+ integer :: d
+ type(c_ptr) :: p
+
+ p = omp_alloc (sizeof(d), ompx_unified_shared_mem_alloc)
+ if (.not. c_associated (p)) stop 1
+
+ do d = 0, omp_get_num_devices ()
+ if (omp_target_is_accessible (p, c_sizeof (d), d) == 0) &
+ stop 2;
+ end do
+
+ call omp_free (p, ompx_unified_shared_mem_alloc);
+end program main