From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id 41E2C384D1A1; Wed, 29 Jun 2022 14:47:32 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 41E2C384D1A1 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] OpenMP, libgomp: Add new runtime routine omp_target_is_accessible. X-Act-Checkin: gcc X-Git-Author: Marcel Vollweiler X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: 67c9b129a857330999c954c6bdf46e8015f489c5 X-Git-Newrev: caa8134568c639dfe980ebab26a514e7e1fa5199 Message-Id: <20220629144732.41E2C384D1A1@sourceware.org> Date: Wed, 29 Jun 2022 14:47:32 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Wed, 29 Jun 2022 14:47:32 -0000 https://gcc.gnu.org/g:caa8134568c639dfe980ebab26a514e7e1fa5199 commit caa8134568c639dfe980ebab26a514e7e1fa5199 Author: Marcel Vollweiler Date: Fri May 6 07:28:26 2022 -0700 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. (cherry picked from commit 4043f53cb4a3541f7a6e4f4132419f78ab7ec4f7) Diff: --- gcc/ChangeLog.omp | 8 ++++ gcc/omp-low.cc | 1 + libgomp/ChangeLog.omp | 15 +++++++ libgomp/libgomp.map | 1 + libgomp/libgomp.texi | 2 +- libgomp/omp.h.in | 2 + libgomp/omp_lib.f90.in | 10 +++++ libgomp/omp_lib.h.in | 11 +++++ libgomp/target.c | 18 ++++++++ .../libgomp.c-c++-common/target-is-accessible-1.c | 47 ++++++++++++++++++++ .../libgomp.fortran/target-is-accessible-1.f90 | 50 ++++++++++++++++++++++ 11 files changed, 164 insertions(+), 1 deletion(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index ac8455b6a3e..59ad9943ae0 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,11 @@ +2022-05-06 Marcel Vollweiler + + Backport from mainline: + 2022-05-06 Marcel Vollweiler + + * omp-low.cc (omp_runtime_api_call): Added target_is_accessible to + omp_runtime_apis array. + 2022-05-05 Sandra Loosemore Backport from mainline: diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 6fa7b31f372..bd7515d6b78 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -4488,6 +4488,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/ChangeLog.omp b/libgomp/ChangeLog.omp index 327b396ca58..d8a7ac2c8ef 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,18 @@ +2022-05-06 Marcel Vollweiler + + Backport from mainline: + 2022-05-06 Marcel Vollweiler + + * 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. + 2022-05-05 Sandra Loosemore Backport from mainline: diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index fc14b8d172b..4423f9c0965 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -229,6 +229,7 @@ OMP_5.1 { OMP_5.1.1 { global: omp_get_mapped_ptr; + omp_target_is_accessible; } OMP_5.1; GOMP_1.0 { diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 317711f3f73..ca959b644d4 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 Y @tab diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in index e811911c99b..5640ad2a70f 100644 --- a/libgomp/omp.h.in +++ b/libgomp/omp.h.in @@ -288,6 +288,8 @@ 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 void *omp_get_mapped_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 d6a251799ac..3b2461819eb 100644 --- a/libgomp/omp_lib.f90.in +++ b/libgomp/omp_lib.f90.in @@ -854,6 +854,16 @@ end function omp_get_mapped_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 0f48510d7ff..28554331c95 100644 --- a/libgomp/omp_lib.h.in +++ b/libgomp/omp_lib.h.in @@ -425,3 +425,14 @@ integer(c_int), value :: device_num end function omp_get_mapped_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 b6298f80962..9a3b44f6e1b 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -3997,6 +3997,24 @@ omp_get_mapped_ptr (const void *ptr, int device_num) return ret; } +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) { 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 00000000000..7c2cf622960 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c @@ -0,0 +1,47 @@ +#include + +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 00000000000..26118553f2d --- /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